[clang] [CIR][AMDGPU] Add support for AMDGCN grid_size builtins (PR #230798)
via cfe-commits
cfe-commits at lists.llvm.org
Sat Oct 10 07:46:27 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Arseniy Obolenskiy (aobolensk)
<details>
<summary>Changes</summary>
Port EmitAMDGPUGridSize from classic codegen
---
Full diff: https://github.com/llvm/llvm-project/pull/230798.diff
4 Files Affected:
- (modified) clang/include/clang/CIR/MissingFeatures.h (+1)
- (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+30-9)
- (modified) clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip (+18)
- (modified) clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip (+20)
``````````diff
diff --git a/clang/include/clang/CIR/MissingFeatures.h b/clang/include/clang/CIR/MissingFeatures.h
index e417be6a89a14..ab5501545d6e7 100644
--- a/clang/include/clang/CIR/MissingFeatures.h
+++ b/clang/include/clang/CIR/MissingFeatures.h
@@ -48,6 +48,7 @@ struct MissingFeatures {
// Load/store attributes
static bool opLoadEmitScalarRangeCheck() { return false; }
+ static bool opLoadRangeNoundefMetadata() { return false; }
static bool opLoadStoreTbaa() { return false; }
static bool opLoadStoreAtomic() { return false; }
static bool opLoadStoreObjC() { return false; }
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 0ee8f18551e41..fa4b87e06e125 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -26,10 +26,9 @@ using namespace cir;
// result to match \p e's return type when needed.
// If \p e is null, returns the raw AS-4 pointer.
static mlir::Value emitAMDGPUDispatchPtr(CIRGenFunction &cgf,
+ mlir::Location loc,
const CallExpr *e = nullptr) {
CIRGenBuilderTy &builder = cgf.getBuilder();
- mlir::Location loc =
- e ? cgf.getLoc(e->getExprLoc()) : builder.getUnknownLoc();
// The intrinsic always returns a pointer in the constant AS.
mlir::Type retTy = cir::PointerType::get(
cir::VoidType::get(builder.getContext()),
@@ -48,6 +47,29 @@ static mlir::Value emitAMDGPUDispatchPtr(CIRGenFunction &cgf,
return builder.createAddrSpaceCast(loc, call, expectedPtrTy);
}
+// 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,
+ mlir::Value offset, mlir::Type ty,
+ CharUnits align) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ mlir::Value addr = builder.createPtrStride(loc, base, offset);
+ cir::LoadOp load = builder.createAlignedLoad(loc, ty, addr, align);
+ load.setInvariant(true);
+ assert(!cir::MissingFeatures::opLoadRangeNoundefMetadata());
+ return load;
+}
+
+static mlir::Value emitAMDGPUGridSize(CIRGenFunction &cgf, mlir::Location loc,
+ unsigned index) {
+ CIRGenBuilderTy &builder = cgf.getBuilder();
+ mlir::Value dispatchPtr = emitAMDGPUDispatchPtr(cgf, loc);
+ // Indexing the HSA kernel_dispatch_packet struct.
+ return emitAMDGPUInvariantLoad(
+ cgf, loc, dispatchPtr, builder.getUInt64(12 + index * 4, loc),
+ builder.getUInt32Ty(), CharUnits::fromQuantity(4));
+}
+
static mlir::Value emitBinaryExpMaybeConstrainedFPBuiltin(
CIRGenFunction &cgf, const CallExpr *e, llvm::StringRef intrinsicName,
llvm::StringRef constrainedIntrinsicName) {
@@ -346,7 +368,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.cos").getValue();
}
case AMDGPU::BI__builtin_amdgcn_dispatch_ptr:
- return emitAMDGPUDispatchPtr(*this, expr);
+ return emitAMDGPUDispatchPtr(*this, getLoc(expr->getExprLoc()), expr);
case AMDGPU::BI__builtin_amdgcn_logf:
case AMDGPU::BI__builtin_amdgcn_log_bf16: {
return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.log").getValue();
@@ -995,14 +1017,13 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
getContext().BuiltinInfo.getName(builtinId));
return mlir::Value{};
}
+ // amdgcn grid size
case AMDGPU::BI__builtin_amdgcn_grid_size_x:
+ return emitAMDGPUGridSize(*this, getLoc(expr->getExprLoc()), 0);
case AMDGPU::BI__builtin_amdgcn_grid_size_y:
- case AMDGPU::BI__builtin_amdgcn_grid_size_z: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- }
+ return emitAMDGPUGridSize(*this, getLoc(expr->getExprLoc()), 1);
+ case AMDGPU::BI__builtin_amdgcn_grid_size_z:
+ return emitAMDGPUGridSize(*this, getLoc(expr->getExprLoc()), 2);
case AMDGPU::BI__builtin_r600_recipsqrt_ieee:
case AMDGPU::BI__builtin_r600_recipsqrt_ieeef: {
cgm.errorNYI(expr->getSourceRange(),
diff --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
index eced98bad1210..29c3b6230031d 100644
--- a/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
+++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-builtins.hip
@@ -84,3 +84,21 @@ __device__ float test_logb(float x) {
// OGCG: load float, ptr addrspace(4) %{{.*}}
// OGCG: call contract{{.*}} @llvm.fabs.f32
// OGCG: fcmp contract one float %{{.*}}, +inf
+
+__device__ unsigned test_grid_size_x() {
+ return __builtin_amdgcn_grid_size_x();
+}
+
+// CIR-LABEL: cir.func no_inline @_Z16test_grid_size_xv
+// CIR: %[[P:.*]] = cir.call_llvm_intrinsic "amdgcn.dispatch.ptr" : () -> !cir.ptr<!void, target_address_space(4)>
+// CIR: cir.load invariant align(4) %{{.*}} : !cir.ptr<!u32i, target_address_space(4)>, !u32i
+
+// LLVM-LABEL: define spir_func noundef i32 @_Z16test_grid_size_xv
+// LLVM: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+// LLVM: [[GEP:%.*]] = getelementptr i8, ptr addrspace(4) [[P]], i64 12
+// LLVM: load i32, ptr addrspace(4) [[GEP]], align 4, !invariant.load
+
+// OGCG-LABEL: define spir_func noundef i32 @_Z16test_grid_size_xv
+// OGCG: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+// OGCG: [[GEP:%.*]] = getelementptr i8, ptr addrspace(4) [[P]], i32 12
+// OGCG: load i32, ptr addrspace(4) [[GEP]], align 4, !range !{{.*}}, !invariant.load
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index 01afd609a9ba9..768936bbcb8d3 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -112,6 +112,26 @@ __device__ void test_dispatch_ptr(__attribute__((address_space(4))) void ** out)
*out = (__attribute__((address_space(4))) void *)__builtin_amdgcn_dispatch_ptr();
}
+// CIR-LABEL: @_Z14test_grid_sizePj
+// CIR: %[[P:.*]] = cir.call_llvm_intrinsic "amdgcn.dispatch.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<!u32i, target_address_space(4)>
+// CIR: cir.load invariant align(4) %[[PTR]]
+// CIR: cir.const #cir.int<16> : !u64i
+// CIR: cir.const #cir.int<20> : !u64i
+// LLVM-LABEL: @_Z14test_grid_sizePj
+// LLVM: [[P:%.*]] = call{{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+// LLVM: [[GEP:%.*]] = getelementptr i8, ptr addrspace(4) [[P]], i{{32|64}} 12
+// LLVM: load i32, ptr addrspace(4) [[GEP]], align 4, {{.*}}!invariant.load
+// LLVM: getelementptr i8, ptr addrspace(4) %{{.*}}, i{{32|64}} 16
+// LLVM: getelementptr i8, ptr addrspace(4) %{{.*}}, i{{32|64}} 20
+__device__ void test_grid_size(unsigned *out) {
+ out[0] = __builtin_amdgcn_grid_size_x();
+ out[1] = __builtin_amdgcn_grid_size_y();
+ out[2] = __builtin_amdgcn_grid_size_z();
+}
+
// CIR-LABEL: @_Z12test_rcp_f32Pff
// CIR: cir.call_llvm_intrinsic "amdgcn.rcp" {{.*}} : (!cir.float) -> !cir.float
// LLVM: define{{.*}} void @_Z12test_rcp_f32Pff
``````````
</details>
https://github.com/llvm/llvm-project/pull/230798
More information about the cfe-commits
mailing list