[clang] [CIR][SPIR-V] Dispatch AMDGPU builtins on AMDGCN-flavored SPIR-V (PR #222018)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Wed Sep 30 04:11:38 PDT 2026
https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/222018
>From ef08c6bd700dec6001269f2d9676a79ea2064d72 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 8 Sep 2026 16:49:31 +0200
Subject: [PATCH 1/3] [CIR][SPIR-V] Dispatch AMDGPU builtins on AMDGCN-flavored
SPIR-V
---
clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 10 +++
clang/lib/CIR/CodeGen/Targets/SPIRV.cpp | 5 ++
.../CIR/CodeGenHIP/amdgcnspirv-builtins.hip | 79 +++++++++++++++++++
3 files changed, 94 insertions(+)
create mode 100644 clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
index d4c17d3c5f24f3..bf64d966fa0bb3 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
@@ -3005,6 +3005,9 @@ RValue CIRGenFunction::emitBuiltinExpr(const GlobalDecl &gd, unsigned builtinID,
llvm::Triple::getArchTypePrefix(getTarget().getTriple().getArch());
if (!prefix.empty()) {
intrinsicID = Intrinsic::getIntrinsicForClangBuiltin(prefix, name);
+ if (intrinsicID == Intrinsic::not_intrinsic && prefix == "spv" &&
+ getTarget().getTriple().getOS() == llvm::Triple::OSType::AMDHSA)
+ intrinsicID = Intrinsic::getIntrinsicForClangBuiltin("amdgcn", name);
// NOTE we don't need to perform a compatibility flag check here since the
// intrinsics are declared in Builtins*.def via LANGBUILTIN which filter the
// MS builtins via ALL_MS_LANGUAGES and are filtered earlier.
@@ -3193,6 +3196,13 @@ emitTargetArchBuiltinExpr(CIRGenFunction *cgf, unsigned builtinID,
case llvm::Triple::riscv32:
case llvm::Triple::riscv64:
return cgf->emitRISCVBuiltinExpr(builtinID, e);
+ case llvm::Triple::spirv32:
+ case llvm::Triple::spirv64:
+ if (cgf->getTarget().getTriple().getOS() == llvm::Triple::OSType::AMDHSA)
+ return cgf->emitAMDGPUBuiltinExpr(builtinID, e);
+ [[fallthrough]];
+ case llvm::Triple::spirv:
+ return std::nullopt;
default:
return std::nullopt;
}
diff --git a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
index 7b66c51af640c6..4d19f122799e3b 100644
--- a/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
+++ b/clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
@@ -43,6 +43,11 @@ class CommonSPIRTargetCIRGenInfo : public TargetCIRGenInfo {
return cir::CallingConv::SpirKernel;
}
+ bool supportsLibCall() const override {
+ const llvm::Triple &triple = getABIInfo().cgt.getCGModule().getTriple();
+ return !(triple.isSPIRV() && triple.getVendor() == llvm::Triple::AMD);
+ }
+
void setCUDAKernelCallingConvention(const FunctionType *&ft) const override {
// Convert HIP kernels to SPIR-V kernels.
if (getABIInfo().cgt.getASTContext().getLangOpts().HIP)
diff --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
new file mode 100644
index 00000000000000..8bcde14d1787cc
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
@@ -0,0 +1,79 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir \
+// RUN: -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir
+
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir \
+// RUN: -fcuda-is-device -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t-cir.ll
+
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip \
+// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=OGCG %s --input-file=%t.ll
+
+// Test that AMDGPU builtins are available on AMDGCN-flavored SPIR-V.
+
+#define __device__ __attribute__((device))
+
+__device__ int test_readfirstlane(int x) {
+ return __builtin_amdgcn_readfirstlane(x);
+}
+
+// CIR-LABEL: cir.func no_inline @_Z18test_readfirstlanei
+// CIR: cir.call_llvm_intrinsic "amdgcn.readfirstlane" {{.*}} : (!s32i) -> !s32i
+
+// LLVM-LABEL: define spir_func noundef i32 @_Z18test_readfirstlanei
+// LLVM: call{{.*}} @llvm.amdgcn.readfirstlane.i32
+
+// OGCG-LABEL: define spir_func noundef i32 @_Z18test_readfirstlanei
+// OGCG: addrspacecast ptr %{{.*}} to ptr addrspace(4)
+// OGCG: call{{.*}} @llvm.amdgcn.readfirstlane.i32
+
+__device__ float test_rcp(float x) {
+ return __builtin_amdgcn_rcpf(x);
+}
+
+// CIR-LABEL: cir.func no_inline @_Z8test_rcpf
+// CIR: cir.call_llvm_intrinsic "amdgcn.rcp" {{.*}} : (!cir.float) -> !cir.float
+
+// LLVM-LABEL: define spir_func noundef float @_Z8test_rcpf
+// LLVM: call{{.*}} @llvm.amdgcn.rcp.f32
+
+// OGCG-LABEL: define spir_func noundef float @_Z8test_rcpf
+// OGCG: call contract{{.*}} @llvm.amdgcn.rcp.f32
+
+// Reached through the generic clang-builtin-to-intrinsic mapping, which has to
+// retry the "amdgcn" prefix after "spv" fails to match.
+
+__device__ unsigned test_wavefrontsize() {
+ return __builtin_amdgcn_wavefrontsize();
+}
+
+// CIR-LABEL: cir.func no_inline @_Z18test_wavefrontsizev
+// CIR: cir.call_llvm_intrinsic "amdgcn.wavefrontsize" : () -> !u32i
+
+// LLVM-LABEL: define spir_func noundef i32 @_Z18test_wavefrontsizev
+// LLVM: call{{.*}} @llvm.amdgcn.wavefrontsize()
+
+// OGCG-LABEL: define spir_func noundef i32 @_Z18test_wavefrontsizev
+// OGCG: call{{.*}} @llvm.amdgcn.wavefrontsize()
+
+// Expanded inline rather than emitted as a libm call, because SPIR-V with an
+// AMD vendor has no device libm.
+
+__device__ float test_logb(float x) {
+ return __builtin_logbf(x);
+}
+
+// CIR-LABEL: cir.func no_inline @_Z9test_logbf
+// CIR: cir.call_llvm_intrinsic "frexp" %{{.*}} : (!cir.float) -> !rec_anon_struct
+
+// LLVM-LABEL: define spir_func noundef float @_Z9test_logbf
+// LLVM: call{{.*}} @llvm.frexp.f32.i32
+// LLVM: add i32 %{{.*}}, -1
+
+// OGCG-LABEL: define spir_func noundef float @_Z9test_logbf
+// OGCG: call{{.*}} @llvm.frexp.f32.i32
+// OGCG: add nsw i32 %{{.*}}, -1
+// OGCG: load float, ptr addrspace(4) %{{.*}}
+// OGCG: call contract{{.*}} @llvm.fabs.f32
>From bc3bb5a2ea831cfddb40cec107182e42419e5acf Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Mon, 28 Sep 2026 20:27:47 +0200
Subject: [PATCH 2/3] address part of the comments
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 5 +++--
clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip | 5 ++++-
2 files changed, 7 insertions(+), 3 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 1c23ea142bf8dc..4c1c4dd9de0f30 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -89,7 +89,8 @@ static mlir::Value emitLogbBuiltin(CIRGenFunction &cgf, const CallExpr *e,
mlir::Value negativeOne =
builder.getConstant(loc, cir::IntAttr::get(int32Ty, -1));
- mlir::Value expMinus1 = builder.createAdd(loc, exp, negativeOne);
+ mlir::Value expMinus1 = builder.createAdd(
+ loc, exp, negativeOne, cir::OverflowBehavior::NoSignedWrap);
mlir::Value siToFp = cir::CastOp::create(
builder, loc, srcTy, cir::CastKind::int_to_float, expMinus1);
@@ -100,7 +101,7 @@ static mlir::Value emitLogbBuiltin(CIRGenFunction &cgf, const CallExpr *e,
mlir::Value inf = builder.getConstant(loc, cir::FPAttr::get(srcTy, infVal));
mlir::Value fabsNegInf =
- builder.createCompare(loc, cir::CmpOpKind::ne, fabs, inf);
+ builder.createCompare(loc, cir::CmpOpKind::one, fabs, inf);
mlir::Value sel = builder.createSelect(loc, fabsNegInf, siToFp, fabs);
diff --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
index 8bcde14d1787cc..36bd58d3f1231a 100644
--- a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
+++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
@@ -70,10 +70,13 @@ __device__ float test_logb(float x) {
// LLVM-LABEL: define spir_func noundef float @_Z9test_logbf
// LLVM: call{{.*}} @llvm.frexp.f32.i32
-// LLVM: add i32 %{{.*}}, -1
+// LLVM: add nsw i32 %{{.*}}, -1
+// LLVM: call{{.*}} @llvm.fabs.f32
+// LLVM: fcmp one float %{{.*}}, +inf
// OGCG-LABEL: define spir_func noundef float @_Z9test_logbf
// OGCG: call{{.*}} @llvm.frexp.f32.i32
// OGCG: add nsw i32 %{{.*}}, -1
// OGCG: load float, ptr addrspace(4) %{{.*}}
// OGCG: call contract{{.*}} @llvm.fabs.f32
+// OGCG: fcmp contract one float %{{.*}}, +inf
>From f34ce4f206666e503e08e5af9c25055cb837d565 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 29 Sep 2026 11:10:50 +0200
Subject: [PATCH 3/3] fixme
---
clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip | 7 +++++++
1 file changed, 7 insertions(+)
diff --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
index 36bd58d3f1231a..3f0e91672d8a5f 100644
--- a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
+++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
@@ -13,6 +13,13 @@
// Test that AMDGPU builtins are available on AMDGCN-flavored SPIR-V.
+// FIXME: CIR doesn't propagate the 'contract' fast-math flag to LLVM IR yet,
+// so the LLVM check lines use call{{.*}} to tolerate the difference between
+// CIR (no flags) and classic codegen ('contract').
+
+// FIXME: CIR doesn't cast allocas to the generic address space yet, so the
+// LLVM output lacks the addrspacecast that classic codegen emits.
+
#define __device__ __attribute__((device))
__device__ int test_readfirstlane(int x) {
More information about the cfe-commits
mailing list