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/3] [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 d4c17d3c5f24f3..bf64d966fa0bb3 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 7b66c51af640c6..4d19f122799e3b 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 00000000000000..8bcde14d1787cc --- /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/3] 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 1c23ea142bf8dc..4c1c4dd9de0f30 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 8bcde14d1787cc..36bd58d3f1231a 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/3] 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 36bd58d3f1231a..3f0e91672d8a5f 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) { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
