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 <[email protected]> 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
