[clang] [CIR][AMDGPU] Add support for AMDGCN workgroup_size builtins (PR #231156)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Mon Oct 12 21:06:40 PDT 2026
https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/231156
>From b567431a78b2883dc6c6a35193d72a42450d62d7 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Mon, 12 Oct 2026 15:16:39 +0200
Subject: [PATCH 1/3] [CIR][AMDGPU] Add support for AMDGCN workgroup_size
builtins
Port EmitAMDGPUWorkGroupSize from classic codegen
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 147 ++++++++++++++++--
.../CIR/CodeGenHIP/amdgcnspirv-builtins.hip | 21 +++
.../CIR/CodeGenHIP/amdgpu-workgroup-size.hip | 57 +++++++
clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 24 +++
4 files changed, 236 insertions(+), 13 deletions(-)
create mode 100644 clang/test/CIR/CodeGenHIP/amdgpu-workgroup-size.hip
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index a9db1082c097f..5e8e61e4b2eba 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -22,6 +22,16 @@ using namespace clang;
using namespace clang::CIRGen;
using namespace cir;
+// Emit a no-operand intrinsic returning a pointer in the constant AS.
+static mlir::Value emitAMDGPUConstantPtrIntrinsic(CIRGenFunction &cgf,
+ mlir::Location loc,
+ llvm::StringRef name) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ mlir::Type retTy = builder.getVoidPtrTy(cir::TargetAddressSpaceAttr::get(
+ builder.getContext(), llvm::AMDGPUAS::CONSTANT_ADDRESS));
+ return builder.emitIntrinsicCallOp(loc, name, retTy, mlir::ValueRange{});
+}
+
// Emit the `amdgcn.dispatch.ptr` intrinsic, address-space-casting the
// result to match \p e's return type when needed.
// If \p e is null, returns the raw AS-4 pointer.
@@ -29,13 +39,8 @@ static mlir::Value emitAMDGPUDispatchPtr(CIRGenFunction &cgf,
mlir::Location loc,
const CallExpr *e = nullptr) {
CIRGenBuilderTy &builder = cgf.getBuilder();
- // The intrinsic always returns a pointer in the constant AS.
- mlir::Type retTy = cir::PointerType::get(
- cir::VoidType::get(builder.getContext()),
- cir::TargetAddressSpaceAttr::get(builder.getContext(),
- llvm::AMDGPUAS::CONSTANT_ADDRESS));
- mlir::Value call = builder.emitIntrinsicCallOp(loc, "amdgcn.dispatch.ptr",
- retTy, mlir::ValueRange{});
+ mlir::Value call =
+ emitAMDGPUConstantPtrIntrinsic(cgf, loc, "amdgcn.dispatch.ptr");
if (!e)
return call;
// Only cast when the caller-visible AS differs from the intrinsic's AS;
@@ -47,6 +52,23 @@ static mlir::Value emitAMDGPUDispatchPtr(CIRGenFunction &cgf,
return builder.createAddrSpaceCast(loc, call, expectedPtrTy);
}
+static mlir::Value emitAMDGPUImplicitArgPtr(CIRGenFunction &cgf,
+ mlir::Location loc) {
+ // OG also marks the result align 8 dereferenceable(256).
+ assert(!cir::MissingFeatures::opCallAttrs());
+ return emitAMDGPUConstantPtrIntrinsic(cgf, loc, "amdgcn.implicitarg.ptr");
+}
+
+static mlir::Value emitAMDGPUWorkGroupID(CIRGenFunction &cgf,
+ mlir::Location loc, unsigned index) {
+ static constexpr llvm::StringLiteral names[] = {"amdgcn.workgroup.id.x",
+ "amdgcn.workgroup.id.y",
+ "amdgcn.workgroup.id.z"};
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ return builder.emitIntrinsicCallOp(loc, names[index], builder.getUInt32Ty(),
+ mlir::ValueRange{});
+}
+
// Load a \p ty value at byte \p offset from \p base as an invariant load.
static mlir::Value emitAMDGPUInvariantLoad(CIRGenFunction &cgf,
mlir::Location loc, mlir::Value base,
@@ -71,6 +93,107 @@ static mlir::Value emitAMDGPUGridSize(CIRGenFunction &cgf, mlir::Location loc,
builder.getUInt32Ty(), CharUnits::fromQuantity(4));
}
+static mlir::Value emitAMDGPUWorkGroupSizeV5(CIRGenFunction &cgf,
+ mlir::Location loc,
+ unsigned index) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ mlir::Value implicitArgPtr = emitAMDGPUImplicitArgPtr(cgf, loc);
+
+ // offsetof(amdhsa_implicit_kernarg_v5, block_count[index])
+ unsigned blockCountOffset = 0 + index * 4;
+ // offsetof(amdhsa_implicit_kernarg_v5, group_size[index])
+ unsigned groupSizeOffset = 12 + index * 2;
+ // offsetof(amdhsa_implicit_kernarg_v5, remainder[index])
+ unsigned remainderOffset = 18 + index * 2;
+
+ mlir::Value sizeOffset;
+ if (cgf.getLangOpts().OffloadUniformBlock) {
+ sizeOffset = builder.getUInt64(groupSizeOffset, loc);
+ } else {
+ mlir::Value blockCount = emitAMDGPUInvariantLoad(
+ cgf, loc, implicitArgPtr, builder.getUInt64(blockCountOffset, loc),
+ builder.getUInt32Ty(), CharUnits::fromQuantity(4));
+ mlir::Value workgroupID = emitAMDGPUWorkGroupID(cgf, loc, index);
+ mlir::Value isFull =
+ builder.createCompare(loc, cir::CmpOpKind::lt, workgroupID, blockCount);
+ sizeOffset = builder.createSelect(loc, isFull,
+ builder.getUInt64(groupSizeOffset, loc),
+ builder.getUInt64(remainderOffset, loc));
+ }
+
+ mlir::Value size = emitAMDGPUInvariantLoad(cgf, loc, implicitArgPtr,
+ sizeOffset, builder.getUInt16Ty(),
+ CharUnits::fromQuantity(2));
+ return builder.createIntCast(size, builder.getUInt32Ty());
+}
+
+static mlir::Value emitAMDGPUWorkGroupSizeV4(CIRGenFunction &cgf,
+ mlir::Location loc,
+ unsigned index) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ mlir::Value dispatchPtr = emitAMDGPUDispatchPtr(cgf, loc);
+
+ // Indexing the HSA kernel_dispatch_packet struct.
+ mlir::Value groupSizeLd = emitAMDGPUInvariantLoad(
+ cgf, loc, dispatchPtr, builder.getUInt64(4 + index * 2, loc),
+ builder.getUInt16Ty(), CharUnits::fromQuantity(2));
+ mlir::Value groupSize =
+ builder.createIntCast(groupSizeLd, builder.getUInt32Ty());
+
+ if (cgf.getLangOpts().OffloadUniformBlock)
+ return groupSize;
+
+ mlir::Value workgroupID = emitAMDGPUWorkGroupID(cgf, loc, index);
+ mlir::Value gridSize = emitAMDGPUInvariantLoad(
+ cgf, loc, dispatchPtr, builder.getUInt64(12 + index * 4, loc),
+ builder.getUInt32Ty(), CharUnits::fromQuantity(4));
+ mlir::Value mul = builder.createMul(loc, workgroupID, groupSize);
+ mlir::Value remainder = builder.createSub(loc, gridSize, mul);
+ mlir::Value isPartial =
+ builder.createCompare(loc, cir::CmpOpKind::lt, remainder, groupSize);
+ return builder.createSelect(loc, isPartial, remainder, groupSize);
+}
+
+static mlir::Value emitAMDGPUWorkGroupSize(CIRGenFunction &cgf,
+ mlir::Location loc, unsigned index) {
+ llvm::CodeObjectVersionKind cov =
+ cgf.getTarget().getTargetOpts().CodeObjectVersion;
+
+ // Do not emit __oclc_ABI_version references with non-empty environment.
+ if (cov == llvm::CodeObjectVersionKind::COV_None &&
+ cgf.getTarget().getTriple().hasEnvironment())
+ cov = llvm::CodeObjectVersionKind::COV_6;
+
+ if (cov != llvm::CodeObjectVersionKind::COV_None)
+ return cov >= llvm::CodeObjectVersionKind::COV_5
+ ? emitAMDGPUWorkGroupSizeV5(cgf, loc, index)
+ : emitAMDGPUWorkGroupSizeV4(cgf, loc, index);
+
+ // Select the implementation at runtime based on __oclc_ABI_version, which
+ // device libraries use to stay ABI-agnostic.
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ CIRGenModule &cgm = cgf.cgm;
+ StringRef name = "__oclc_ABI_version";
+ mlir::Type int32Ty = builder.getSInt32Ty();
+ auto abiVersionGV =
+ mlir::dyn_cast_or_null<cir::GlobalOp>(cgm.getGlobalValue(name));
+ if (!abiVersionGV)
+ abiVersionGV =
+ cgm.createGlobalOp(loc, name, int32Ty, /*isConstant=*/false,
+ cir::toCIRAddressSpaceAttr(cgm.getMLIRContext(),
+ LangAS::opencl_constant));
+ mlir::Value abiVersionPtr = builder.createGetGlobal(loc, abiVersionGV);
+ mlir::Value abiVersion = builder.createAlignedLoad(
+ loc, int32Ty, abiVersionPtr,
+ cgf.getContext().getTypeAlignInChars(cgf.getContext().IntTy));
+ mlir::Value isCOV5 = builder.createCompare(
+ loc, cir::CmpOpKind::ge, abiVersion,
+ builder.getConstantInt(loc, int32Ty, llvm::CodeObjectVersionKind::COV_5));
+ mlir::Value v5Impl = emitAMDGPUWorkGroupSizeV5(cgf, loc, index);
+ mlir::Value v4Impl = emitAMDGPUWorkGroupSizeV4(cgf, loc, index);
+ return builder.createSelect(loc, isCOV5, v5Impl, v4Impl);
+}
+
static mlir::Value emitBinaryExpMaybeConstrainedFPBuiltin(
CIRGenFunction &cgf, const CallExpr *e, llvm::StringRef intrinsicName,
llvm::StringRef constrainedIntrinsicName) {
@@ -1011,13 +1134,11 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
}
// amdgcn workgroup size
case AMDGPU::BI__builtin_amdgcn_workgroup_size_x:
+ return emitAMDGPUWorkGroupSize(*this, getLoc(expr->getExprLoc()), 0);
case AMDGPU::BI__builtin_amdgcn_workgroup_size_y:
- case AMDGPU::BI__builtin_amdgcn_workgroup_size_z: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- }
+ return emitAMDGPUWorkGroupSize(*this, getLoc(expr->getExprLoc()), 1);
+ case AMDGPU::BI__builtin_amdgcn_workgroup_size_z:
+ return emitAMDGPUWorkGroupSize(*this, getLoc(expr->getExprLoc()), 2);
// amdgcn grid size
case AMDGPU::BI__builtin_amdgcn_grid_size_x:
return emitAMDGPUGridSize(*this, getLoc(expr->getExprLoc()), 0);
diff --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
index 25aa1c4a7d187..e2c30de872562 100644
--- a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
+++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
@@ -138,3 +138,24 @@ __device__ unsigned test_grid_size_z() {
// OGCG: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
// OGCG: [[GEP:%.*]] = getelementptr i8, ptr addrspace(4) [[P]], i32 20
// OGCG: load i32, ptr addrspace(4) [[GEP]], align 4, !range !{{.*}}, !invariant.load
+
+__device__ unsigned test_workgroup_size_x() {
+ return __builtin_amdgcn_workgroup_size_x();
+}
+
+// CIR-LABEL: cir.func no_inline @_Z21test_workgroup_size_xv
+// CIR: %[[P:.*]] = cir.call_llvm_intrinsic "amdgcn.implicitarg.ptr" : () -> !cir.ptr<!void, target_address_space(4)>
+// CIR: %[[GEP:.*]] = cir.ptr_stride %[[P]], %{{.*}} : (!cir.ptr<!void, target_address_space(4)>, !u64i)
+// CIR: cir.load invariant align(2) %{{.*}} : !cir.ptr<!u16i, target_address_space(4)>, !u16i
+
+// LLVM-LABEL: define spir_func noundef i32 @_Z21test_workgroup_size_xv
+// LLVM: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+// LLVM: [[GEP:%.*]] = getelementptr i8, ptr addrspace(4) [[P]], i64 12
+// LLVM: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, !invariant.load
+// LLVM: zext i16 [[SIZE]] to i32
+
+// OGCG-LABEL: define spir_func noundef i32 @_Z21test_workgroup_size_xv
+// OGCG: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+// OGCG: [[GEP:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[P]], i64 12
+// OGCG: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, !range !{{.*}}, !invariant.load !{{.*}}, !noundef
+// OGCG: zext i16 [[SIZE]] to i32
diff --git a/clang/test/CIR/CodeGenHIP/amdgpu-workgroup-size.hip b/clang/test/CIR/CodeGenHIP/amdgpu-workgroup-size.hip
new file mode 100644
index 0000000000000..e5b3f28efc3be
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/amdgpu-workgroup-size.hip
@@ -0,0 +1,57 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fclangir -fcuda-is-device \
+// RUN: -mcode-object-version=4 -fno-offload-uniform-block -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=COV4 %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -mcode-object-version=4 -fno-offload-uniform-block -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=COV4 %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fclangir -fcuda-is-device \
+// RUN: -mcode-object-version=5 -fno-offload-uniform-block -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=COV5 %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -mcode-object-version=5 -fno-offload-uniform-block -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=COV5 %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fclangir -fcuda-is-device \
+// RUN: -mcode-object-version=none -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=NONE %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device \
+// RUN: -mcode-object-version=none -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=NONE %s
+
+#define __device__ __attribute__((device))
+
+// NONE: @__oclc_ABI_version = external addrspace(4) global i32
+
+// COV4-LABEL: @_Z4testv
+// COV4: [[DP:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+// COV4: [[GEP:%.*]] = getelementptr{{.*}} i8, ptr addrspace(4) [[DP]], i64 6
+// COV4: [[LD:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2
+// COV4: [[GS:%.*]] = zext i16 [[LD]] to i32
+// COV4: [[ID:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()
+// COV4: [[GGEP:%.*]] = getelementptr{{.*}} i8, ptr addrspace(4) [[DP]], i64 16
+// COV4: [[GRID:%.*]] = load i32, ptr addrspace(4) [[GGEP]], align 4
+// COV4: [[MUL:%.*]] = mul i32 [[ID]], [[GS]]
+// COV4: [[REM:%.*]] = sub i32 [[GRID]], [[MUL]]
+// COV4: [[CMP:%.*]] = icmp ult i32 [[REM]], [[GS]]
+// COV4: select i1 [[CMP]], i32 [[REM]], i32 [[GS]]
+
+// COV5-LABEL: @_Z4testv
+// COV5: [[IP:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+// COV5: [[BGEP:%.*]] = getelementptr{{.*}} i8, ptr addrspace(4) [[IP]], i64 4
+// COV5: [[BC:%.*]] = load i32, ptr addrspace(4) [[BGEP]], align 4
+// COV5: [[ID:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()
+// COV5: [[FULL:%.*]] = icmp ult i32 [[ID]], [[BC]]
+// COV5: [[OFF:%.*]] = select i1 [[FULL]], i{{32|64}} 14, i{{32|64}} 20
+// COV5: [[SGEP:%.*]] = getelementptr{{.*}} i8, ptr addrspace(4) [[IP]], i{{32|64}} [[OFF]]
+// COV5: [[SZ:%.*]] = load i16, ptr addrspace(4) [[SGEP]], align 2
+// COV5: zext i16 [[SZ]] to i32
+
+// NONE-LABEL: @_Z4testv
+// NONE: [[ABI:%.*]] = load i32, ptr addrspace(4) @__oclc_ABI_version, align 4
+// NONE: [[IS5:%.*]] = icmp sge i32 [[ABI]], 500
+// NONE: call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+// NONE: [[V5:%.*]] = zext i16 %{{.*}} to i32
+// NONE: call{{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+// NONE: [[V4:%.*]] = zext i16 %{{.*}} to i32
+// NONE: select i1 [[IS5]], i32 [[V5]], i32 [[V4]]
+__device__ unsigned test() { return __builtin_amdgcn_workgroup_size_y(); }
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index 768936bbcb8d3..180f632b5050b 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -132,6 +132,30 @@ __device__ void test_grid_size(unsigned *out) {
out[2] = __builtin_amdgcn_grid_size_z();
}
+// CIR-LABEL: @_Z19test_workgroup_sizePj
+// CIR: %[[P:.*]] = cir.call_llvm_intrinsic "amdgcn.implicitarg.ptr" : () -> !cir.ptr<!void, target_address_space(4)>
+// CIR: %[[OFF:.*]] = cir.const #cir.int<12> : !u64i
+// CIR: %[[GEP:.*]] = cir.ptr_stride %[[P]], %[[OFF]]
+// CIR: %[[PTR:.*]] = cir.cast bitcast %[[GEP]] : !cir.ptr<!void, target_address_space(4)> -> !cir.ptr<!u16i, target_address_space(4)>
+// CIR: %[[SIZE:.*]] = cir.load invariant align(2) %[[PTR]]
+// CIR: cir.cast integral %[[SIZE]] : !u16i -> !u32i
+// CIR: cir.const #cir.int<14> : !u64i
+// CIR: cir.const #cir.int<16> : !u64i
+// LLVM-LABEL: @_Z19test_workgroup_sizePj
+// LLVM: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+// LLVM: [[GEP:%.*]] = getelementptr{{.*}} i8, ptr addrspace(4) [[P]], i64 12
+// LLVM: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, {{.*}}!invariant.load
+// LLVM: zext i16 [[SIZE]] to i32
+// LLVM: getelementptr{{.*}} i8, ptr addrspace(4) %{{.*}}, i64 14
+// LLVM: getelementptr{{.*}} i8, ptr addrspace(4) %{{.*}}, i64 16
+// LLVM-NOT: @llvm.amdgcn.workgroup.id
+// LLVM: ret void
+__device__ void test_workgroup_size(unsigned *out) {
+ out[0] = __builtin_amdgcn_workgroup_size_x();
+ out[1] = __builtin_amdgcn_workgroup_size_y();
+ out[2] = __builtin_amdgcn_workgroup_size_z();
+}
+
// CIR-LABEL: @_Z12test_rcp_f32Pff
// CIR: cir.call_llvm_intrinsic "amdgcn.rcp" {{.*}} : (!cir.float) -> !cir.float
// LLVM: define{{.*}} void @_Z12test_rcp_f32Pff
>From 926396f53b01281abd48b72361dcfb5e0595ed91 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 13 Oct 2026 06:05:49 +0200
Subject: [PATCH 2/3] Address comments
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 4 +-
.../CIR/CodeGenHIP/amdgcnspirv-builtins.hip | 42 +++++++++++++++++++
2 files changed, 44 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 5e8e61e4b2eba..a0cb98beacd85 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -154,6 +154,8 @@ static mlir::Value emitAMDGPUWorkGroupSizeV4(CIRGenFunction &cgf,
return builder.createSelect(loc, isPartial, remainder, groupSize);
}
+// COV_NONE picks COV_4 or COV_5+ at runtime from __oclc_ABI_version so device
+// libraries stay ABI-agnostic.
static mlir::Value emitAMDGPUWorkGroupSize(CIRGenFunction &cgf,
mlir::Location loc, unsigned index) {
llvm::CodeObjectVersionKind cov =
@@ -169,8 +171,6 @@ static mlir::Value emitAMDGPUWorkGroupSize(CIRGenFunction &cgf,
? emitAMDGPUWorkGroupSizeV5(cgf, loc, index)
: emitAMDGPUWorkGroupSizeV4(cgf, loc, index);
- // Select the implementation at runtime based on __oclc_ABI_version, which
- // device libraries use to stay ABI-agnostic.
CIRGenBuilderTy &builder = cgf.getBuilder();
CIRGenModule &cgm = cgf.cgm;
StringRef name = "__oclc_ABI_version";
diff --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
index e2c30de872562..99e2ba1c8e237 100644
--- a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
+++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
@@ -159,3 +159,45 @@ __device__ unsigned test_workgroup_size_x() {
// OGCG: [[GEP:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[P]], i64 12
// OGCG: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, !range !{{.*}}, !invariant.load !{{.*}}, !noundef
// OGCG: zext i16 [[SIZE]] to i32
+
+__device__ unsigned test_workgroup_size_y() {
+ return __builtin_amdgcn_workgroup_size_y();
+}
+
+// CIR-LABEL: cir.func no_inline @_Z21test_workgroup_size_yv
+// CIR: %[[P:.*]] = cir.call_llvm_intrinsic "amdgcn.implicitarg.ptr" : () -> !cir.ptr<!void, target_address_space(4)>
+// CIR: %[[GEP:.*]] = cir.ptr_stride %[[P]], %{{.*}} : (!cir.ptr<!void, target_address_space(4)>, !u64i)
+// CIR: cir.load invariant align(2) %{{.*}} : !cir.ptr<!u16i, target_address_space(4)>, !u16i
+
+// LLVM-LABEL: define spir_func noundef i32 @_Z21test_workgroup_size_yv
+// LLVM: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+// LLVM: [[GEP:%.*]] = getelementptr i8, ptr addrspace(4) [[P]], i64 14
+// LLVM: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, !invariant.load
+// LLVM: zext i16 [[SIZE]] to i32
+
+// OGCG-LABEL: define spir_func noundef i32 @_Z21test_workgroup_size_yv
+// OGCG: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+// OGCG: [[GEP:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[P]], i64 14
+// OGCG: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, !range !{{.*}}, !invariant.load !{{.*}}, !noundef
+// OGCG: zext i16 [[SIZE]] to i32
+
+__device__ unsigned test_workgroup_size_z() {
+ return __builtin_amdgcn_workgroup_size_z();
+}
+
+// CIR-LABEL: cir.func no_inline @_Z21test_workgroup_size_zv
+// CIR: %[[P:.*]] = cir.call_llvm_intrinsic "amdgcn.implicitarg.ptr" : () -> !cir.ptr<!void, target_address_space(4)>
+// CIR: %[[GEP:.*]] = cir.ptr_stride %[[P]], %{{.*}} : (!cir.ptr<!void, target_address_space(4)>, !u64i)
+// CIR: cir.load invariant align(2) %{{.*}} : !cir.ptr<!u16i, target_address_space(4)>, !u16i
+
+// LLVM-LABEL: define spir_func noundef i32 @_Z21test_workgroup_size_zv
+// LLVM: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+// LLVM: [[GEP:%.*]] = getelementptr i8, ptr addrspace(4) [[P]], i64 16
+// LLVM: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, !invariant.load
+// LLVM: zext i16 [[SIZE]] to i32
+
+// OGCG-LABEL: define spir_func noundef i32 @_Z21test_workgroup_size_zv
+// OGCG: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+// OGCG: [[GEP:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[P]], i64 16
+// OGCG: [[SIZE:%.*]] = load i16, ptr addrspace(4) [[GEP]], align 2, !range !{{.*}}, !invariant.load !{{.*}}, !noundef
+// OGCG: zext i16 [[SIZE]] to i32
>From 3252278b2cb5a1fc554b67f9f605fe7558837ff6 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 13 Oct 2026 06:06:29 +0200
Subject: [PATCH 3/3] orig comment
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 9 +++++++--
1 file changed, 7 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index a0cb98beacd85..3a0cdcce7872d 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -154,8 +154,13 @@ static mlir::Value emitAMDGPUWorkGroupSizeV4(CIRGenFunction &cgf,
return builder.createSelect(loc, isPartial, remainder, groupSize);
}
-// COV_NONE picks COV_4 or COV_5+ at runtime from __oclc_ABI_version so device
-// libraries stay ABI-agnostic.
+// \p index is 0, 1, and 2 for x, y, and z dimension, respectively.
+// Emit code based on Code Object ABI version.
+// COV_4 : Emit code to use dispatch ptr
+// COV_5+ : Emit code to use implicitarg ptr
+// COV_NONE : Emit code to load a global variable "__oclc_ABI_version"
+// and use its value for COV_4 or COV_5+ approach. It is used for
+// compiling device libraries in an ABI-agnostic way.
static mlir::Value emitAMDGPUWorkGroupSize(CIRGenFunction &cgf,
mlir::Location loc, unsigned index) {
llvm::CodeObjectVersionKind cov =
More information about the cfe-commits
mailing list