llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir Author: Chaitanya (skc7) <details> <summary>Changes</summary> Port `emitAMDGPUDispatchPtr` from OGCG. Emits the `amdgcn.dispatch.ptr` intrinsic and inserts an address-space cast when the builtin's expected return type differs. --- Full diff: https://github.com/llvm/llvm-project/pull/199880.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+29-6) - (modified) clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip (+8) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 7e6e0f1a06046..8e35b24f1e501 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -14,12 +14,39 @@ #include "mlir/IR/Value.h" #include "clang/Basic/TargetBuiltins.h" +#include "llvm/Support/AMDGPUAddrSpace.h" #include "llvm/Support/ErrorHandling.h" 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. +// 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(); + // 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{}); + if (!e) + return call; + // Only cast when the caller-visible AS differs from the intrinsic's AS; + auto expectedPtrTy = + mlir::cast<cir::PointerType>(cgf.convertType(e->getType())); + auto callPtrTy = mlir::cast<cir::PointerType>(call.getType()); + if (expectedPtrTy.getAddrSpace() == callPtrTy.getAddrSpace()) + return call; + return builder.createAddrSpaceCast(loc, call, expectedPtrTy); +} + static mlir::Value emitBinaryExpMaybeConstrainedFPBuiltin( CIRGenFunction &cgf, const CallExpr *e, llvm::StringRef intrinsicName, llvm::StringRef constrainedIntrinsicName) { @@ -268,12 +295,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(); +} `````````` </details> https://github.com/llvm/llvm-project/pull/199880 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
