[clang] a3b8700 - [CIR][SPIR-V] Dispatch AMDGPU builtins on AMDGCN-flavored SPIR-V (#222018)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 2 02:48:56 PDT 2026
Author: Arseniy Obolenskiy
Date: 2026-10-02T09:48:48Z
New Revision: a3b8700c9576ca89753a0cc23423f37ef3b6c17d
URL: https://github.com/llvm/llvm-project/commit/a3b8700c9576ca89753a0cc23423f37ef3b6c17d
DIFF: https://github.com/llvm/llvm-project/commit/a3b8700c9576ca89753a0cc23423f37ef3b6c17d.diff
LOG: [CIR][SPIR-V] Dispatch AMDGPU builtins on AMDGCN-flavored SPIR-V (#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
Added:
clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
Modified:
clang/include/clang/CIR/MissingFeatures.h
clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
clang/lib/CIR/CodeGen/Targets/SPIRV.cpp
clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/CIR/MissingFeatures.h b/clang/include/clang/CIR/MissingFeatures.h
index 9a866b62849ef..bbf76c976c7e8 100644
--- a/clang/include/clang/CIR/MissingFeatures.h
+++ b/clang/include/clang/CIR/MissingFeatures.h
@@ -26,6 +26,7 @@ namespace cir {
struct MissingFeatures {
// Address space related
static bool addressSpace() { return false; }
+ static bool spirvDefaultIsGenericAddrSpace() { return false; }
// Unhandled global/linkage information.
static bool opGlobalThreadLocal() { return false; }
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
index fe7a63d7bb85d..b090c876397e4 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
@@ -3173,6 +3173,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.
@@ -3361,6 +3364,11 @@ 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);
+ 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/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp
index 5367b4c76e2a0..084d98f81bce2 100644
--- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp
@@ -8,6 +8,7 @@
#include "../TargetLoweringInfo.h"
#include "clang/CIR/Dialect/IR/CIROpsEnums.h"
+#include "clang/CIR/MissingFeatures.h"
namespace cir {
@@ -29,6 +30,9 @@ class SPIRVTargetLoweringInfo : public TargetLoweringInfo {
public:
unsigned getTargetAddrSpaceFromCIRAddrSpace(
cir::LangAddressSpace addrSpace) const override {
+ // TODO(cir): SYCL and CUDA/HIP device code map Default to Generic
+ // (SPIRDefIsGenMap).
+ assert(!cir::MissingFeatures::spirvDefaultIsGenericAddrSpace());
auto idx = static_cast<unsigned>(addrSpace);
assert(idx < std::size(SPIRVAddrSpaceMap) &&
"Unknown CIR address space for SPIR-V target");
diff --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
new file mode 100644
index 0000000000000..015d4e57a2e29
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
@@ -0,0 +1,92 @@
+// 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.
+
+// FIXME: CIR doesn't propagate the 'contract' fast-math flag to LLVM IR yet.
+// The LLVM checks match the flag-free form and fail once it lands.
+
+// FIXME: CIR lowers the Default address space to 0 rather than generic on
+// SPIR-V, so the LLVM output lacks the addrspacecast that classic codegen
+// emits. The LLVM checks fail once it lands.
+
+#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-NOT: addrspacecast
+// 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 addrspace(4) float @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
+
+// FIXME: CIR emits 'add' without 'nsw' and 'fcmp une' instead of 'fcmp one'.
+
+// LLVM-LABEL: define spir_func noundef float @_Z9test_logbf
+// LLVM: call{{.*}} @llvm.frexp.f32.i32
+// LLVM: add i32 %{{.*}}, -1
+// LLVM: call addrspace(4) float @llvm.fabs.f32
+// LLVM: fcmp une 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
More information about the cfe-commits
mailing list