Author: Chaitanya Date: 2026-03-12T10:20:02+05:30 New Revision: 65d0ef49e136ecd1cac3b3a9d5cb3a4a0d134396
URL: https://github.com/llvm/llvm-project/commit/65d0ef49e136ecd1cac3b3a9d5cb3a4a0d134396 DIFF: https://github.com/llvm/llvm-project/commit/65d0ef49e136ecd1cac3b3a9d5cb3a4a0d134396.diff LOG: [CIR][AMDGPU] Add AMDGPU target support to CIR CodeGen (#185819) This PR adds AMDGPUTargetCIRGenInfo and AMDGPUABIInfo to handle the amdgcn triple in CIR code generation, along with a basic HIP codegen test. Added: clang/test/CIR/CodeGenHIP/simple.cpp Modified: clang/lib/CIR/CodeGen/CIRGenModule.cpp clang/lib/CIR/CodeGen/TargetInfo.cpp clang/lib/CIR/CodeGen/TargetInfo.h Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 25fa4c7c86a89..cb931f969a41d 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -266,6 +266,10 @@ const TargetCIRGenInfo &CIRGenModule::getTargetCIRGenInfo() { case llvm::Triple::nvptx64: theTargetCIRGenInfo = createNVPTXTargetCIRGenInfo(genTypes); return *theTargetCIRGenInfo; + case llvm::Triple::amdgcn: { + theTargetCIRGenInfo = createAMDGPUTargetCIRGenInfo(genTypes); + return *theTargetCIRGenInfo; + } } } diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp index 70ffb46050ea1..ee68d9c329b83 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.cpp +++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp @@ -44,6 +44,21 @@ bool clang::CIRGen::isEmptyFieldForLayout(const ASTContext &context, namespace { +class AMDGPUABIInfo : public ABIInfo { +public: + AMDGPUABIInfo(CIRGenTypes &cgt) : ABIInfo(cgt) {} +}; + +class AMDGPUTargetCIRGenInfo : public TargetCIRGenInfo { +public: + AMDGPUTargetCIRGenInfo(CIRGenTypes &cgt) + : TargetCIRGenInfo(std::make_unique<AMDGPUABIInfo>(cgt)) {} +}; + +} // namespace + +namespace { + class X8664ABIInfo : public ABIInfo { public: X8664ABIInfo(CIRGenTypes &cgt) : ABIInfo(cgt) {} @@ -71,6 +86,11 @@ class NVPTXTargetCIRGenInfo : public TargetCIRGenInfo { }; } // namespace +std::unique_ptr<TargetCIRGenInfo> +clang::CIRGen::createAMDGPUTargetCIRGenInfo(CIRGenTypes &cgt) { + return std::make_unique<AMDGPUTargetCIRGenInfo>(cgt); +} + std::unique_ptr<TargetCIRGenInfo> clang::CIRGen::createNVPTXTargetCIRGenInfo(CIRGenTypes &cgt) { return std::make_unique<NVPTXTargetCIRGenInfo>(cgt); diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h index 9ba155b220fbc..b397d8cd7fab8 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.h +++ b/clang/lib/CIR/CodeGen/TargetInfo.h @@ -132,6 +132,9 @@ class TargetCIRGenInfo { } }; +std::unique_ptr<TargetCIRGenInfo> +createAMDGPUTargetCIRGenInfo(CIRGenTypes &cgt); + std::unique_ptr<TargetCIRGenInfo> createX8664TargetCIRGenInfo(CIRGenTypes &cgt); std::unique_ptr<TargetCIRGenInfo> createNVPTXTargetCIRGenInfo(CIRGenTypes &cgt); diff --git a/clang/test/CIR/CodeGenHIP/simple.cpp b/clang/test/CIR/CodeGenHIP/simple.cpp new file mode 100644 index 0000000000000..bc046816309ab --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/simple.cpp @@ -0,0 +1,89 @@ +#include "../CodeGenCUDA/Inputs/cuda.h" + +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir \ +// RUN: -x hip -fhip-new-launch-api \ +// RUN: -I%S/../CodeGenCUDA/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%S/../CodeGenCUDA/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 \ +// RUN: -x hip -emit-llvm -fhip-new-launch-api \ +// RUN: -I%S/../CodeGenCUDA/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: -fvisibility=hidden \ +// RUN: -I%S/../CodeGenCUDA/Inputs/ -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=OGCG-DEVICE --input-file=%t.ll %s + +// Per Thread Stream test cases: + +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \ +// RUN: -fgpu-default-stream=per-thread -DHIP_API_PER_THREAD_DEFAULT_STREAM \ +// RUN: -x hip -emit-llvm -fhip-new-launch-api \ +// RUN: -I%S/../CodeGenCUDA/Inputs/ %s -o %t.ll +// RUN: FileCheck --check-prefix=OGCG-HOST-PTH --input-file=%t.ll %s + + +__host__ void host_fn(int *a, int *b, int *c) {} +// CIR-HOST: cir.func {{.*}} @_Z7host_fnPiS_S_ +// CIR-DEVICE-NOT: cir.func {{.*}} @_Z7host_fnPiS_S_ + +__device__ void device_fn(int* a, double b, float c) {} +// CIR-HOST-NOT: cir.func {{.*}} @_Z9device_fnPidf +// CIR-DEVICE: cir.func {{.*}} @_Z9device_fnPidf + +__global__ void global_fn(int a) {} +// CIR-DEVICE: cir.func {{.*}} @_Z9global_fni +// OGCG-DEVICE: define protected amdgpu_kernel void @_Z9global_fni + +// CIR-HOST: @_Z24__device_stub__global_fni{{.*}}attributes {cu.kernel_name = #cir.cu.kernel_name<_Z9global_fni>} +// CIR-HOST: %[[#CIRKernelArgs:]] = cir.alloca {{.*}}"kernel_args" +// CIR-HOST: %[[#Decayed:]] = cir.cast array_to_ptrdecay %[[#CIRKernelArgs]] +// CIR-HOST: cir.call @__hipPopCallConfiguration +// CIR-HOST: cir.get_global @_Z9global_fni : !cir.ptr<!cir.func<(!s32i)>> +// CIR-HOST: cir.call @hipLaunchKernel + +// 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 +// OGCG-HOST-PTH: %call = call noundef i32 @hipLaunchKernel_spt(ptr noundef @_Z9global_fni + + +int main() { + global_fn<<<1, 1>>>(1); +} +// CIR-DEVICE-NOT: cir.func {{.*}} @main() + +// CIR-HOST: cir.func {{.*}} @main() +// CIR-HOST: cir.call @_ZN4dim3C1Ejjj +// CIR-HOST: cir.call @_ZN4dim3C1Ejjj +// CIR-HOST: [[Push:%[0-9]+]] = cir.call @__hipPushCallConfiguration +// CIR-HOST: [[ConfigOK:%[0-9]+]] = cir.cast int_to_bool [[Push]] +// CIR-HOST: cir.if [[ConfigOK]] { +// CIR-HOST: } else { +// CIR-HOST: [[Arg:%[0-9]+]] = cir.const #cir.int<1> +// CIR-HOST: cir.call @_Z24__device_stub__global_fni([[Arg]]) +// CIR-HOST: } + +// 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
