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 <[email protected]> Date: Tue, 8 Sep 2026 16:49:31 +0200 Subject: [PATCH 1/4] [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 >From bc3bb5a2ea831cfddb40cec107182e42419e5acf Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Mon, 28 Sep 2026 20:27:47 +0200 Subject: [PATCH 2/4] 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 1c23ea142bf8d..4c1c4dd9de0f3 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 8bcde14d1787c..36bd58d3f1231 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 <[email protected]> Date: Tue, 29 Sep 2026 11:10:50 +0200 Subject: [PATCH 3/4] 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 36bd58d3f1231..3f0e91672d8a5 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) { >From 3f66ce32c2a25acc94265fa973835ec07c3353ce Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Thu, 1 Oct 2026 10:41:33 +0200 Subject: [PATCH 4/4] comments --- clang/include/clang/CIR/MissingFeatures.h | 1 + clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 2 -- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 5 ++--- .../TargetLowering/Targets/SPIRV.cpp | 4 ++++ .../CIR/CodeGenHIP/amdgcnspirv-builtins.hip | 21 +++++++++++-------- 5 files changed, 19 insertions(+), 14 deletions(-) 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 88c843f52f009..b090c876397e4 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -3368,8 +3368,6 @@ emitTargetArchBuiltinExpr(CIRGenFunction *cgf, unsigned builtinID, 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/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 02175f6ec4110..b2aac0d376fa1 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -89,8 +89,7 @@ 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, cir::OverflowBehavior::NoSignedWrap); + mlir::Value expMinus1 = builder.createAdd(loc, exp, negativeOne); mlir::Value siToFp = cir::CastOp::create( builder, loc, srcTy, cir::CastKind::int_to_float, expMinus1); @@ -101,7 +100,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::one, fabs, inf); + builder.createCompare(loc, cir::CmpOpKind::ne, fabs, inf); mlir::Value sel = builder.createSelect(loc, fabsNegInf, siToFp, fabs); 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 index 3f0e91672d8a5..015d4e57a2e29 100644 --- a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip +++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip @@ -13,12 +13,12 @@ // 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 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 doesn't cast allocas to the generic address space yet, so the -// LLVM output lacks the addrspacecast that classic codegen emits. +// 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)) @@ -30,6 +30,7 @@ __device__ int test_readfirstlane(int x) { // 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 @@ -44,7 +45,7 @@ __device__ float test_rcp(float x) { // 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 +// 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 @@ -75,11 +76,13 @@ __device__ float test_logb(float 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 nsw i32 %{{.*}}, -1 -// LLVM: call{{.*}} @llvm.fabs.f32 -// LLVM: fcmp one float %{{.*}}, +inf +// 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
