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