[CIR][AMDGPU] Add AMDGPU target support to CIR CodeGen - #185819
Merged
Merged
Conversation
skc7
marked this pull request as ready for review
March 11, 2026 06:59
skc7
requested review from
andykaylor,
bcardosolopes,
lanza and
xlauko
as code owners
March 11, 2026 06:59
llvmbot
added
clang
Member
|
@llvm/pr-subscribers-clang @llvm/pr-subscribers-clangir Author: Chaitanya (skc7) ChangesThis PR adds AMDGPUTargetCIRGenInfo and AMDGPUABIInfo to handle the amdgcn triple in CIR code generation, along with a Note: changes related to amdgpu have been pulled from #179084 to unblock amdgpu clangIR PRs upstreaming. Full diff: https://github.com/llvm/llvm-project/pull/185819.diff 4 Files Affected:
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
|
skc7
force-pushed
the
users/skc7/clangir/amdgpu_target
branch
from
March 11, 2026 09:33
5e43deb to
032b2ff
Compare
RiverDave
approved these changes
Mar 12, 2026
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/66/builds/27757 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/186/builds/16987 Here is the relevant piece of the build log for the reference |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
5 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
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 #179084 to unblock amdgpu clangIR PRs upstreaming.