https://github.com/Kalomidin updated https://github.com/llvm/llvm-project/pull/193701
>From d1450e28ce7d2e37450ba321cbe1e017a4a3f7d4 Mon Sep 17 00:00:00 2001 From: Kalomidin <[email protected]> Date: Thu, 14 May 2026 09:36:33 +0800 Subject: [PATCH] [NVPTX] Allow mixed address-space pointer arguments in kernel lowering --- .../CodeGenCUDA/multi-pointer-kernel-args.cu | 9 +++++++ .../Target/NVPTX/NVPTXTargetTransformInfo.cpp | 25 +++++++++++++++++++ .../Target/NVPTX/NVPTXTargetTransformInfo.h | 1 + 3 files changed, 35 insertions(+) create mode 100644 clang/test/CodeGenCUDA/multi-pointer-kernel-args.cu diff --git a/clang/test/CodeGenCUDA/multi-pointer-kernel-args.cu b/clang/test/CodeGenCUDA/multi-pointer-kernel-args.cu new file mode 100644 index 0000000000000..b8c59fcc29c02 --- /dev/null +++ b/clang/test/CodeGenCUDA/multi-pointer-kernel-args.cu @@ -0,0 +1,9 @@ +// RUN: %clang_cc1 -x cuda --cuda-device-only %s -S -o - | FileCheck %s + +// CHECK: st.global.b32 +// CHECK: st.global.b32 + +__global__ void kernel(int **X, int x, int y) { + X[x][y] = x * y; + X[y][x] = x + y; +} diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp index a491d0ed4a912..1c7ae1b58b742 100644 --- a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.cpp @@ -652,7 +652,32 @@ unsigned NVPTXTTIImpl::getAssumedAddrSpace(const Value *V) const { return ADDRESS_SPACE_LOCAL; } } + if (int AS = getPointerLoadAddressSpace(V); AS != -1) { + return AS; + } + return -1; +} +/* + * geparg = getelementptr T, ptr %arg, i32 0, i32 0 + * ptr a = load ptr %geparg + * we can expect AS of a to be global + */ +int NVPTXTTIImpl::getPointerLoadAddressSpace(const Value *V) const { + auto *Load = dyn_cast<LoadInst>(V); + if (!Load) + return -1; + auto *Ptr = Load->getPointerOperand(); + if (!Ptr) + return -1; + // if it is argument, return GM AS + if (const auto *Arg = dyn_cast<Argument>(Ptr)) { + if (isKernelFunction(*Arg->getParent())) + return ADDRESS_SPACE_GLOBAL; + } + if (auto *GEP = dyn_cast<GetElementPtrInst>(Ptr)) { + return getAssumedAddrSpace(GEP->getPointerOperand()); + } return -1; } diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h index 8bdafd6b905f1..87a816e939f44 100644 --- a/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h +++ b/llvm/lib/Target/NVPTX/NVPTXTargetTransformInfo.h @@ -212,6 +212,7 @@ class NVPTXTTIImpl final : public BasicTTIImplBase<NVPTXTTIImpl> { Value *rewriteIntrinsicWithAddressSpace(IntrinsicInst *II, Value *OldV, Value *NewV) const override; unsigned getAssumedAddrSpace(const Value *V) const override; + int getPointerLoadAddressSpace(const Value *V) const; void collectKernelLaunchBounds( const Function &F, _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
