From 3f7bc449d9aecbce1063b1737b5bbc2654578cbc Mon Sep 17 00:00:00 2001 From: koparasy Date: Tue, 4 Nov 2025 15:21:40 -0800 Subject: [PATCH 1/5] Initial lowering to LLVM-IR for device code --- .../Transforms/TargetLowering/CMakeLists.txt | 1 + .../Transforms/TargetLowering/LowerModule.cpp | 2 + .../Transforms/TargetLowering/TargetInfo.h | 3 + .../TargetLowering/Targets/AMDGPU.cpp | 72 +++++++++++++++++++ .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 33 ++++++--- clang/test/CIR/CodeGen/HIP/simple.cpp | 7 +- 6 files changed, 106 insertions(+), 12 deletions(-) create mode 100644 clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt index dab8dbbe5611..c4fdab93fe2f 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CMakeLists.txt @@ -13,6 +13,7 @@ add_clang_library(TargetLowering TargetInfo.cpp TargetLoweringInfo.cpp Targets/AArch64.cpp + Targets/AMDGPU.cpp Targets/NVPTX.cpp Targets/SPIR.cpp Targets/X86.cpp diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp index d040c87282f5..6578e9df8b42 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/LowerModule.cpp @@ -71,6 +71,8 @@ createTargetLoweringInfo(LowerModule &LM) { return createAArch64TargetLoweringInfo(LM, Kind); } + case llvm::Triple::amdgcn: + return createAMDGPUTargetLoweringInfo(LM); case llvm::Triple::x86_64: { switch (Triple.getOS()) { case llvm::Triple::Win32: diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetInfo.h b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetInfo.h index a03cf711babc..3a6ba03a2aa2 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetInfo.h +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/TargetInfo.h @@ -33,6 +33,9 @@ createSPIRVTargetLoweringInfo(LowerModule &CGM); std::unique_ptr createNVPTXTargetLoweringInfo(LowerModule &CGM); +std::unique_ptr +createAMDGPUTargetLoweringInfo(LowerModule &CGM); + } // namespace cir #endif // LLVM_CLANG_LIB_CIR_DIALECT_TRANSFORMS_TARGETLOWERING_TARGETINFO_H diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp new file mode 100644 index 000000000000..1af2b12a8195 --- /dev/null +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp @@ -0,0 +1,72 @@ +//===- AMDGPU.cpp - TargetInfo for AMDGPU +//-----------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "ABIInfoImpl.h" +#include "LowerFunctionInfo.h" +#include "LowerTypes.h" +#include "TargetInfo.h" +#include "TargetLoweringInfo.h" +#include "clang/CIR/ABIArgInfo.h" +#include "clang/CIR/Dialect/IR/CIRTypes.h" +#include "clang/CIR/MissingFeatures.h" +#include "llvm/Support/Casting.h" +#include "llvm/Support/ErrorHandling.h" + +using ABIArgInfo = cir::ABIArgInfo; +using MissingFeature = cir::MissingFeatures; + +namespace cir { + +//===----------------------------------------------------------------------===// +// AMDGPU ABI Implementation +//===----------------------------------------------------------------------===// + +namespace { + +class AMDGPUABIInfo : public ABIInfo { +public: + AMDGPUABIInfo(LowerTypes <) : ABIInfo(lt) {} + +private: + void computeInfo(LowerFunctionInfo &fi) const override { + llvm_unreachable("NYI"); + } +}; + +class AMDGPUTargetLoweringInfo : public TargetLoweringInfo { +public: + AMDGPUTargetLoweringInfo(LowerTypes <) + : TargetLoweringInfo(std::make_unique(lt)) {} + // Taken from here: https://llvm.org/docs/AMDGPUUsage.html#address-spaces + unsigned getTargetAddrSpaceFromCIRAddrSpace( + cir::AddressSpace addrSpace) const override { + switch (addrSpace) { + case cir::AddressSpace::OffloadPrivate: + return 5; + case cir::AddressSpace::OffloadLocal: + return 3; + case cir::AddressSpace::OffloadGlobal: + return 1; + case cir::AddressSpace::OffloadConstant: + return 4; + case cir::AddressSpace::OffloadGeneric: + return 0; + default: + cir_cconv_unreachable("Unknown CIR address space for this target"); + } + } +}; + +} // namespace +std::unique_ptr +createAMDGPUTargetLoweringInfo(LowerModule &lowerModule) { + return std::make_unique(lowerModule.getTypes()); +} + +} // namespace cir diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index 9308a278e107..7b1c7a2ec2ca 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -1749,19 +1749,31 @@ mlir::LogicalResult CIRToLLVMAllocaOpLowering::matchAndRewrite( convertTypeForMemory(*getTypeConverter(), dataLayout, op.getAllocaType()); auto resultTy = getTypeConverter()->convertType(op.getType()); // Verification between the CIR alloca AS and the one from data layout. - { + auto allocaAS = [&]() { auto resPtrTy = mlir::cast(resultTy); auto dlAllocaASAttr = mlir::cast_if_present( dataLayout.getAllocaMemorySpace()); - // Absence means 0 // TODO: The query for the alloca AS should be done through CIRDataLayout // instead to reuse the logic of interpret null attr as 0. - auto dlAllocaAS = dlAllocaASAttr ? dlAllocaASAttr.getInt() : 0; - if (dlAllocaAS != resPtrTy.getAddressSpace()) { - return op.emitError() << "alloca address space doesn't match the one " - "from the target data layout: " - << dlAllocaAS; - } + if (!dlAllocaASAttr) + return 0u; + return static_cast(dlAllocaASAttr.getValue().getZExtValue()); + }(); + + auto resPtrTy = + mlir::LLVM::LLVMPointerType::get(elementTy.getContext(), allocaAS); + + auto llvmAlloca = rewriter.create( + op.getLoc(), resPtrTy, elementTy, size, op.getAlignmentAttr().getInt()); + + auto expectedPtrTy = mlir::cast( + getTypeConverter()->convertType(op.getResult().getType())); + + mlir::Value finalPtr = llvmAlloca.getResult(); + + if (expectedPtrTy.getAddressSpace() != allocaAS) { + finalPtr = rewriter.create( + op.getLoc(), expectedPtrTy, finalPtr); } // If there are annotations available, copy them out before we destroy the @@ -1770,11 +1782,10 @@ mlir::LogicalResult CIRToLLVMAllocaOpLowering::matchAndRewrite( if (op.getAnnotations()) annotations = op.getAnnotationsAttr(); - auto llvmAlloca = rewriter.replaceOpWithNewOp( - op, resultTy, elementTy, size, op.getAlignmentAttr().getInt()); - if (annotations && !annotations.empty()) buildAllocaAnnotations(llvmAlloca, adaptor, rewriter, annotations); + rewriter.replaceOp(op, finalPtr); + return mlir::success(); } diff --git a/clang/test/CIR/CodeGen/HIP/simple.cpp b/clang/test/CIR/CodeGen/HIP/simple.cpp index aaed4e31e192..7bd333e29e92 100644 --- a/clang/test/CIR/CodeGen/HIP/simple.cpp +++ b/clang/test/CIR/CodeGen/HIP/simple.cpp @@ -15,6 +15,11 @@ // RUN: %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM-HOST --input-file=%t.ll %s +// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \ +// RUN: -fcuda-is-device -fhip-new-launch-api \ +// RUN: -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM-DEVICE --input-file=%t.ll %s + // Attribute for global_fn // CIR-HOST: [[Kernel:#[a-zA-Z_0-9]+]] = {{.*}}#cir.cu.kernel_name<_Z9global_fni>{{.*}} @@ -29,6 +34,7 @@ __device__ void device_fn(int* a, double b, float c) {} __global__ void global_fn(int a) {} // CIR-DEVICE: @_Z9global_fni +// LLVM-DEVICE: define dso_local void @_Z9global_fni // CIR-HOST: @_Z24__device_stub__global_fni{{.*}}extra([[Kernel]]) // CIR-HOST: %[[#CIRKernelArgs:]] = cir.alloca {{.*}}"kernel_args" @@ -44,7 +50,6 @@ __global__ void global_fn(int a) {} // LLVM-HOST: call i32 @__hipPopCallConfiguration // LLVM-HOST: call i32 @hipLaunchKernel(ptr @_Z9global_fni - int main() { global_fn<<<1, 1>>>(1); } From f5e8acc9b3d2bd9b402e4c3d3275d7d5cd34c233 Mon Sep 17 00:00:00 2001 From: koparasy Date: Wed, 5 Nov 2025 15:25:55 -0800 Subject: [PATCH 2/5] Add addr-space test --- .../CIR/CodeGen/HIP/addrspace-lowering.cpp | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) create mode 100644 clang/test/CIR/CodeGen/HIP/addrspace-lowering.cpp diff --git a/clang/test/CIR/CodeGen/HIP/addrspace-lowering.cpp b/clang/test/CIR/CodeGen/HIP/addrspace-lowering.cpp new file mode 100644 index 000000000000..5c72c35186d8 --- /dev/null +++ b/clang/test/CIR/CodeGen/HIP/addrspace-lowering.cpp @@ -0,0 +1,19 @@ +#include "../Inputs/cuda.h" + +// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \ +// RUN: -fcuda-is-device -fhip-new-launch-api \ +// RUN: -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM-DEVICE --input-file=%t.ll %s + + +__shared__ int a; + +// LLVM-DEVICE: @a = addrspace(3) {{.*}} + +__device__ int b; + +// LLVM-DEVICE: @b = addrspace(1) {{.*}} + +__constant__ int c; + +// LLVM-DEVICE: @c = addrspace(4) {{.*}} From 3f2b24de9c9e660291d8743974ce99008ac7197b Mon Sep 17 00:00:00 2001 From: koparasy Date: Wed, 5 Nov 2025 15:35:11 -0800 Subject: [PATCH 3/5] Check CIR codegen of address spaces --- clang/test/CIR/CodeGen/HIP/address-spaces.cpp | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) create mode 100644 clang/test/CIR/CodeGen/HIP/address-spaces.cpp diff --git a/clang/test/CIR/CodeGen/HIP/address-spaces.cpp b/clang/test/CIR/CodeGen/HIP/address-spaces.cpp new file mode 100644 index 000000000000..86c5d8d90918 --- /dev/null +++ b/clang/test/CIR/CodeGen/HIP/address-spaces.cpp @@ -0,0 +1,19 @@ +#include "../Inputs/cuda.h" + +// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \ +// RUN: -fcuda-is-device -fhip-new-launch-api \ +// RUN: -emit-cir %s -o %t.ll +// RUN: FileCheck --check-prefix=CIR --input-file=%t.ll %s + +__global__ void fn() { + int i = 0; + __shared__ int j; + j = i; +} + +// CIR: cir.global "private" internal dso_local addrspace(offload_local) @_ZZ2fnvE1j : !s32i +// CIR: cir.func dso_local @_Z2fnv +// CIR: [[Local:%[0-9]+]] = cir.alloca !s32i, !cir.ptr, ["i", init] +// CIR: [[Shared:%[0-9]+]] = cir.get_global @_ZZ2fnvE1j : !cir.ptr +// CIR: [[Tmp:%[0-9]+]] = cir.load {{.*}} [[Local]] : !cir.ptr, !s32i +// CIR: cir.store{{.*}} [[Tmp]], [[Shared]] : !s32i, !cir.ptr From 7cad52383599069b71bf004ef7ba8679d87838ff Mon Sep 17 00:00:00 2001 From: koparasy Date: Thu, 6 Nov 2025 22:07:14 -0800 Subject: [PATCH 4/5] Fix comments --- .../TargetLowering/Targets/AMDGPU.cpp | 3 +- clang/test/CIR/CodeGen/HIP/address-spaces.cpp | 4 +- .../CIR/CodeGen/HIP/addrspace-lowering.cpp | 20 ++++++--- clang/test/CIR/CodeGen/HIP/simple.cpp | 44 ++++++++++++++++--- 4 files changed, 55 insertions(+), 16 deletions(-) diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp index 1af2b12a8195..7432972889ed 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp @@ -1,5 +1,4 @@ -//===- AMDGPU.cpp - TargetInfo for AMDGPU -//-----------------------------------===// +//===- AMDGPU.cpp - TargetInfo for AMDGPU ---------------------------------===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. diff --git a/clang/test/CIR/CodeGen/HIP/address-spaces.cpp b/clang/test/CIR/CodeGen/HIP/address-spaces.cpp index 86c5d8d90918..1b028c16fe2d 100644 --- a/clang/test/CIR/CodeGen/HIP/address-spaces.cpp +++ b/clang/test/CIR/CodeGen/HIP/address-spaces.cpp @@ -1,8 +1,8 @@ -#include "../Inputs/cuda.h" +#include "cuda.h" // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \ // RUN: -fcuda-is-device -fhip-new-launch-api \ -// RUN: -emit-cir %s -o %t.ll +// RUN: -I$(dirname %s)/../Inputs/ -emit-cir %s -o %t.ll // RUN: FileCheck --check-prefix=CIR --input-file=%t.ll %s __global__ void fn() { diff --git a/clang/test/CIR/CodeGen/HIP/addrspace-lowering.cpp b/clang/test/CIR/CodeGen/HIP/addrspace-lowering.cpp index 5c72c35186d8..40e771e2af1d 100644 --- a/clang/test/CIR/CodeGen/HIP/addrspace-lowering.cpp +++ b/clang/test/CIR/CodeGen/HIP/addrspace-lowering.cpp @@ -1,19 +1,25 @@ -#include "../Inputs/cuda.h" +#include "cuda.h" // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \ // RUN: -fcuda-is-device -fhip-new-launch-api \ -// RUN: -emit-llvm %s -o %t.ll +// RUN: -I$(dirname %s)/../Inputs/ -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM-DEVICE --input-file=%t.ll %s +// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \ +// RUN: -fcuda-is-device -fhip-new-launch-api \ +// RUN: -I$(dirname %s)/../Inputs/ -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=OGCG-DEVICE --input-file=%t.ll %s -__shared__ int a; -// LLVM-DEVICE: @a = addrspace(3) {{.*}} +__shared__ int a; +// LLVM-DEVICE: @a = addrspace(3) global i32 undef, align 4 +// OGCG-DEVICE: @a = addrspace(3) global i32 undef, align 4 __device__ int b; - -// LLVM-DEVICE: @b = addrspace(1) {{.*}} +// LLVM-DEVICE: @b = addrspace(1) externally_initialized global i32 0, align 4 +// OGCG-DEVICE: @b = addrspace(1) externally_initialized global i32 0, align 4 __constant__ int c; +// LLVM-DEVICE: @c = addrspace(4) externally_initialized constant i32 0, align 4 +// OGCG-DEVICE: @c = addrspace(4) externally_initialized constant i32 0, align 4 -// LLVM-DEVICE: @c = addrspace(4) {{.*}} diff --git a/clang/test/CIR/CodeGen/HIP/simple.cpp b/clang/test/CIR/CodeGen/HIP/simple.cpp index 7bd333e29e92..3a04b8697ae5 100644 --- a/clang/test/CIR/CodeGen/HIP/simple.cpp +++ b/clang/test/CIR/CodeGen/HIP/simple.cpp @@ -1,25 +1,36 @@ -#include "../Inputs/cuda.h" +#include "cuda.h" // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \ // RUN: -x hip -fhip-new-launch-api \ -// RUN: -emit-cir %s -o %t.cir +// RUN: -I$(dirname %s)/../Inputs/ -emit-cir %s -o %t.cir // RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \ // RUN: -fcuda-is-device -fhip-new-launch-api \ -// RUN: -emit-cir %s -o %t.cir +// RUN: -I$(dirname %s)/../Inputs/ -emit-cir %s -o %t.cir // RUN: FileCheck --check-prefix=CIR-DEVICE --input-file=%t.cir %s // // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \ // RUN: -x hip -emit-llvm -fhip-new-launch-api \ -// RUN: %s -o %t.ll +// RUN: -I$(dirname %s)/../Inputs/ %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM-HOST --input-file=%t.ll %s // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \ // RUN: -fcuda-is-device -fhip-new-launch-api \ -// RUN: -emit-llvm %s -o %t.ll +// RUN: -I$(dirname %s)/../Inputs/ -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM-DEVICE --input-file=%t.ll %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \ +// RUN: -x hip -emit-llvm -fhip-new-launch-api \ +// RUN: -I$(dirname %s)/../Inputs/ %s -o %t.ll +// RUN: FileCheck --check-prefix=OGCG-HOST --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \ +// RUN: -fcuda-is-device -fhip-new-launch-api \ +// RUN: -I$(dirname %s)/../Inputs/ -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=OGCG-DEVICE --input-file=%t.ll %s + + // Attribute for global_fn // CIR-HOST: [[Kernel:#[a-zA-Z_0-9]+]] = {{.*}}#cir.cu.kernel_name<_Z9global_fni>{{.*}} @@ -35,6 +46,7 @@ __device__ void device_fn(int* a, double b, float c) {} __global__ void global_fn(int a) {} // CIR-DEVICE: @_Z9global_fni // LLVM-DEVICE: define dso_local void @_Z9global_fni +// OGCG-DEVICE: define dso_local amdgpu_kernel void @_Z9global_fni // CIR-HOST: @_Z24__device_stub__global_fni{{.*}}extra([[Kernel]]) // CIR-HOST: %[[#CIRKernelArgs:]] = cir.alloca {{.*}}"kernel_args" @@ -49,6 +61,13 @@ __global__ void global_fn(int a) {} // LLVM-HOST: %[[#GEP2:]] = getelementptr [1 x ptr], ptr %[[#KernelArgs]], i32 0, i64 0 // LLVM-HOST: call i32 @__hipPopCallConfiguration // LLVM-HOST: call i32 @hipLaunchKernel(ptr @_Z9global_fni +// +// OGCG-HOST: define dso_local void @_Z24__device_stub__global_fni +// OGCG-HOST: %kernel_args = alloca ptr, i64 1, align 16 +// OGCG-HOST: getelementptr ptr, ptr %kernel_args, i32 0 +// OGCG-HOST: call i32 @__hipPopCallConfiguration +// OGCG-HOST: %call = call noundef i32 @hipLaunchKernel(ptr noundef @_Z9global_fni + int main() { global_fn<<<1, 1>>>(1); @@ -83,3 +102,18 @@ int main() { // LLVM-HOST: %[[#]] = load i32 // LLVM-HOST: ret i32 +// OGCG-HOST: define dso_local noundef i32 @main +// OGCG-HOST: %agg.tmp = alloca %struct.dim3, align 4 +// OGCG-HOST: %agg.tmp1 = alloca %struct.dim3, align 4 +// OGCG-HOST: call void @_ZN4dim3C1Ejjj +// OGCG-HOST: call void @_ZN4dim3C1Ejjj +// OGCG-HOST: %call = call i32 @__hipPushCallConfiguration +// OGCG-HOST: %tobool = icmp ne i32 %call, 0 +// OGCG-HOST: br i1 %tobool, label %kcall.end, label %kcall.configok +// OGCG-HOST: kcall.configok: +// OGCG-HOST: call void @_Z24__device_stub__global_fni(i32 noundef 1) +// OGCG-HOST: br label %kcall.end +// OGCG-HOST: kcall.end: +// OGCG-HOST: %{{[0-9]+}} = load i32, ptr %retval, align 4 +// OGCG-HOST: ret i32 %8 + From 223ec78f45def53cb2d68abad80b4e3b28655eb2 Mon Sep 17 00:00:00 2001 From: koparasy Date: Fri, 7 Nov 2025 15:55:43 -0800 Subject: [PATCH 5/5] Use %S to get path --- clang/test/CIR/CodeGen/HIP/address-spaces.cpp | 2 +- clang/test/CIR/CodeGen/HIP/addrspace-lowering.cpp | 4 ++-- clang/test/CIR/CodeGen/HIP/simple.cpp | 12 ++++++------ 3 files changed, 9 insertions(+), 9 deletions(-) diff --git a/clang/test/CIR/CodeGen/HIP/address-spaces.cpp b/clang/test/CIR/CodeGen/HIP/address-spaces.cpp index 1b028c16fe2d..3ac0c30e1fe1 100644 --- a/clang/test/CIR/CodeGen/HIP/address-spaces.cpp +++ b/clang/test/CIR/CodeGen/HIP/address-spaces.cpp @@ -2,7 +2,7 @@ // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \ // RUN: -fcuda-is-device -fhip-new-launch-api \ -// RUN: -I$(dirname %s)/../Inputs/ -emit-cir %s -o %t.ll +// RUN: -I%S/../Inputs/ -emit-cir %s -o %t.ll // RUN: FileCheck --check-prefix=CIR --input-file=%t.ll %s __global__ void fn() { diff --git a/clang/test/CIR/CodeGen/HIP/addrspace-lowering.cpp b/clang/test/CIR/CodeGen/HIP/addrspace-lowering.cpp index 40e771e2af1d..e45810e39bf1 100644 --- a/clang/test/CIR/CodeGen/HIP/addrspace-lowering.cpp +++ b/clang/test/CIR/CodeGen/HIP/addrspace-lowering.cpp @@ -2,12 +2,12 @@ // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \ // RUN: -fcuda-is-device -fhip-new-launch-api \ -// RUN: -I$(dirname %s)/../Inputs/ -emit-llvm %s -o %t.ll +// RUN: -I%S/../Inputs/ -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM-DEVICE --input-file=%t.ll %s // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \ // RUN: -fcuda-is-device -fhip-new-launch-api \ -// RUN: -I$(dirname %s)/../Inputs/ -emit-llvm %s -o %t.ll +// RUN: -I%S/../Inputs/ -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=OGCG-DEVICE --input-file=%t.ll %s diff --git a/clang/test/CIR/CodeGen/HIP/simple.cpp b/clang/test/CIR/CodeGen/HIP/simple.cpp index 3a04b8697ae5..1f1049856d8d 100644 --- a/clang/test/CIR/CodeGen/HIP/simple.cpp +++ b/clang/test/CIR/CodeGen/HIP/simple.cpp @@ -2,32 +2,32 @@ // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \ // RUN: -x hip -fhip-new-launch-api \ -// RUN: -I$(dirname %s)/../Inputs/ -emit-cir %s -o %t.cir +// RUN: -I%S/../Inputs/ -emit-cir %s -o %t.cir // RUN: FileCheck --check-prefix=CIR-HOST --input-file=%t.cir %s // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \ // RUN: -fcuda-is-device -fhip-new-launch-api \ -// RUN: -I$(dirname %s)/../Inputs/ -emit-cir %s -o %t.cir +// RUN: -I%S/../Inputs/ -emit-cir %s -o %t.cir // RUN: FileCheck --check-prefix=CIR-DEVICE --input-file=%t.cir %s // // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \ // RUN: -x hip -emit-llvm -fhip-new-launch-api \ -// RUN: -I$(dirname %s)/../Inputs/ %s -o %t.ll +// RUN: -I%S/../Inputs/ %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM-HOST --input-file=%t.ll %s // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip -fclangir \ // RUN: -fcuda-is-device -fhip-new-launch-api \ -// RUN: -I$(dirname %s)/../Inputs/ -emit-llvm %s -o %t.ll +// RUN: -I%S/../Inputs/ -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM-DEVICE --input-file=%t.ll %s // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \ // RUN: -x hip -emit-llvm -fhip-new-launch-api \ -// RUN: -I$(dirname %s)/../Inputs/ %s -o %t.ll +// RUN: -I%S/../Inputs/ %s -o %t.ll // RUN: FileCheck --check-prefix=OGCG-HOST --input-file=%t.ll %s // RUN: %clang_cc1 -triple=amdgcn-amd-amdhsa -x hip \ // RUN: -fcuda-is-device -fhip-new-launch-api \ -// RUN: -I$(dirname %s)/../Inputs/ -emit-llvm %s -o %t.ll +// RUN: -I%S/../Inputs/ -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=OGCG-DEVICE --input-file=%t.ll %s