[clang] [CIR][AMDGPU] Add AMDGPU target support to CIR CodeGen (PR #185819)
via cfe-commits
cfe-commits at lists.llvm.org
Tue Mar 10 23:58:00 PDT 2026
https://github.com/skc7 created https://github.com/llvm/llvm-project/pull/185819
This PR adds AMDGPUTargetCIRGenInfo and AMDGPUABIInfo to handle the amdgcn triple in CIR code generation, along with a
basic HIP codegen test.
Note: changes related to amdgpu have been pulled from https://github.com/llvm/llvm-project/pull/179084 to unblock amdgpu clangIR PRs upstreaming.
>From 5e43debd2b5aa992fcc2b9038d4b34bb2b683857 Mon Sep 17 00:00:00 2001
From: skc7 <Krishna.Sankisa at amd.com>
Date: Wed, 11 Mar 2026 12:14:11 +0530
Subject: [PATCH] [CIR][AMDGPU] Add AMDGPU target support to CIR CodeGen
---
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 5 +-
clang/lib/CIR/CodeGen/TargetInfo.cpp | 20 ++++++
clang/lib/CIR/CodeGen/TargetInfo.h | 3 +
clang/test/CIR/CodeGenHIP/simple.cpp | 89 ++++++++++++++++++++++++++
4 files changed, 116 insertions(+), 1 deletion(-)
create mode 100644 clang/test/CIR/CodeGenHIP/simple.cpp
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 1b4939b4741bb..12897ab8fe7a8 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -241,7 +241,10 @@ const TargetCIRGenInfo &CIRGenModule::getTargetCIRGenInfo() {
// Currently we just fall through to x86_64.
[[fallthrough]];
-
+ case llvm::Triple::amdgcn: {
+ theTargetCIRGenInfo = createAMDGPUTargetCIRGenInfo(genTypes);
+ return *theTargetCIRGenInfo;
+ }
case llvm::Triple::x86_64: {
switch (triple.getOS()) {
default:
diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp
index 2f3824d3d47a7..b32b383777817 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 f4792d5309e36..dad36740639b4 100644
--- a/clang/lib/CIR/CodeGen/TargetInfo.h
+++ b/clang/lib/CIR/CodeGen/TargetInfo.h
@@ -116,6 +116,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
More information about the cfe-commits
mailing list