[clang] [CIR][SPIR-V] Dispatch AMDGPU builtins on AMDGCN-flavored SPIR-V (PR #222018)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 1 01:41:50 PDT 2026
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 <arseniy.obolenskiy at amd.com>
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 <arseniy.obolenskiy at amd.com>
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 <arseniy.obolenskiy at amd.com>
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 <arseniy.obolenskiy at amd.com>
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
More information about the cfe-commits
mailing list