https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/230798
Port EmitAMDGPUGridSize from classic codegen >From d5e5b449494de4da717b39a61fb09f8605e0b46c Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Sat, 10 Oct 2026 16:43:48 +0200 Subject: [PATCH] [CIR][AMDGPU] Add support for AMDGCN grid_size builtins --- clang/include/clang/CIR/MissingFeatures.h | 1 + clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 39 ++++++++++++++----- .../CIR/CodeGenHIP/amdgcnspirv-builtins.hip | 18 +++++++++ clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 20 ++++++++++ 4 files changed, 69 insertions(+), 9 deletions(-) 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
