https://github.com/skc7 created https://github.com/llvm/llvm-project/pull/199880
None >From 034ba70b2e835efde2af91dbe8d9ede31be704af Mon Sep 17 00:00:00 2001 From: skc7 <[email protected]> Date: Wed, 27 May 2026 11:24:55 +0530 Subject: [PATCH] [CIR][AMDGPU] Implement lowering for __builtin_amdgcn_dispatch_ptr --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 29 +++++++++++++++---- clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 8 +++++ 2 files changed, 31 insertions(+), 6 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 7e6e0f1a06046..00e427e4d7fe3 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -20,6 +20,27 @@ using namespace clang; using namespace clang::CIRGen; using namespace cir; +// Emit the `amdgcn.dispatch.ptr` intrinsic, address-space-casting the +// result to match \p e's return type when needed (e.g. HIP flat AS). +// If \p e is null, returns the raw AS-4 pointer. +static mlir::Value emitAMDGPUDispatchPtr(CIRGenFunction &cgf, + const CallExpr *e = nullptr) { + CIRGenBuilderTy &builder = cgf.getBuilder(); + mlir::Location loc = e ? cgf.getLoc(e->getExprLoc()) + : builder.getUnknownLoc(); + mlir::Type retTy = cir::PointerType::get( + cir::VoidType::get(builder.getContext()), + cir::TargetAddressSpaceAttr::get(builder.getContext(), 4)); + mlir::Value call = builder.emitIntrinsicCallOp(loc, "amdgcn.dispatch.ptr", + retTy, mlir::ValueRange{}); + if (!e) + return call; + mlir::Type expectedTy = cgf.convertType(e->getType()); + if (expectedTy == call.getType()) + return call; + return builder.createAddrSpaceCast(loc, call, expectedTy); +} + static mlir::Value emitBinaryExpMaybeConstrainedFPBuiltin( CIRGenFunction &cgf, const CallExpr *e, llvm::StringRef intrinsicName, llvm::StringRef constrainedIntrinsicName) { @@ -268,12 +289,8 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, getContext().BuiltinInfo.getName(builtinId)); return mlir::Value{}; } - case AMDGPU::BI__builtin_amdgcn_dispatch_ptr: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } + case AMDGPU::BI__builtin_amdgcn_dispatch_ptr: + return emitAMDGPUDispatchPtr(*this, expr); case AMDGPU::BI__builtin_amdgcn_logf: case AMDGPU::BI__builtin_amdgcn_log_bf16: { cgm.errorNYI(expr->getSourceRange(), diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index b3c5ac59679a3..08ae9f52cb379 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -87,3 +87,11 @@ __device__ void test_readlane(int* out, int a, int b) { __device__ void test_readfirstlane(int* out, int a) { *out = __builtin_amdgcn_readfirstlane(a); } + +// CIR-LABEL: @_Z17test_dispatch_ptr +// CIR: %{{.*}} = cir.call_llvm_intrinsic "amdgcn.dispatch.ptr" : () -> !cir.ptr<!void, target_address_space(4)> +// LLVM-LABEL: @_Z17test_dispatch_ptr +// LLVM: call{{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr() +__device__ void test_dispatch_ptr(__attribute__((address_space(4))) void ** out) { + *out = (__attribute__((address_space(4))) void *)__builtin_amdgcn_dispatch_ptr(); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
