[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
Tue Sep 8 07:52:33 PDT 2026
https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/222018
Mappings to classic codegen:
- spirv32/spirv64 arch switch: https://github.com/llvm/llvm-project/blob/52d922aa153b1232813470e014635419ce006023/clang/lib/CodeGen/CGBuiltin.cpp#L129-L135
- spv -> amdgcn prefix handling: https://github.com/llvm/llvm-project/blob/52d922aa153b1232813470e014635419ce006023/clang/lib/CodeGen/CGBuiltin.cpp#L6879-L6881
- `supportsLibCall()`: https://github.com/llvm/llvm-project/blob/52d922aa153b1232813470e014635419ce006023/clang/lib/CodeGen/Targets/SPIR.cpp#L139-L142
>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] [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 d4c17d3c5f24f..bf64d966fa0bb 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 7b66c51af640c..4d19f122799e3 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 0000000000000..8bcde14d1787c
--- /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
More information about the cfe-commits
mailing list