[llvm] [NVPTX] Lower allocas to the local address space (PR #204346)
Tim Besard via llvm-commits
llvm-commits at lists.llvm.org
Wed Jul 22 22:48:27 PDT 2026
https://github.com/maleadt updated https://github.com/llvm/llvm-project/pull/204346
>From 21bcec84549f0944501fdbeed6861e17214c2091 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Fri, 5 Jun 2026 08:31:03 +0200
Subject: [PATCH 01/15] Replace generic allocas with local allocas
NVPTXLowerAlloca used to insert generic-to-local-to-generic addrspacecast pairs so NVPTXInferAddressSpaces could recover local stack accesses. Instead, replace generic allocas with addrspace(5) allocas and rewrite localizable users directly to the local pointer.
Escaping users keep a single local-to-generic cast. This exposes stack memory to later lowering without teaching generic SelectionDAG code to look through addrspacecasts.
---
llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp | 275 +++++++++++++-----
llvm/test/CodeGen/NVPTX/lower-alloca.ll | 15 +-
llvm/test/CodeGen/NVPTX/lower-byval-args.ll | 133 +++++----
.../test/CodeGen/NVPTX/memcpy-alloca-align.ll | 132 +++++++++
llvm/test/CodeGen/NVPTX/variadics-backend.ll | 107 +++----
5 files changed, 465 insertions(+), 197 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/memcpy-alloca-align.ll
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
index ef8df4d70d956..51d8b1ab3b9b1 100644
--- a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
@@ -6,27 +6,31 @@
//
//===----------------------------------------------------------------------===//
//
-// For all alloca instructions, and add a pair of cast to local address for
-// each of them. For example,
+// For generic alloca instructions, create an equivalent alloca in local
+// address space and rewrite uses that can directly use local memory. For
+// example,
//
// %A = alloca i32
// store i32 0, i32* %A ; emits st.u32
//
// will be transformed to
//
-// %A = alloca i32
-// %Local = addrspacecast i32* %A to i32 addrspace(5)*
-// %Generic = addrspacecast i32 addrspace(5)* %A to i32*
-// store i32 0, i32 addrspace(5)* %Generic ; emits st.local.u32
+// %A = alloca i32, addrspace(5)
+// store i32 0, i32 addrspace(5)* %A ; emits st.local.u32
//
-// And we will rely on NVPTXInferAddressSpaces to combine the last two
-// instructions.
+// Uses that require a generic pointer use a single addrspacecast from the
+// local alloca.
//
//===----------------------------------------------------------------------===//
#include "MCTargetDesc/NVPTXBaseInfo.h"
#include "NVPTX.h"
+#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/SmallPtrSet.h"
+#include "llvm/ADT/SmallVector.h"
#include "llvm/IR/Function.h"
+#include "llvm/IR/IntrinsicInst.h"
+#include "llvm/IR/Intrinsics.h"
#include "llvm/IR/Instructions.h"
#include "llvm/IR/Type.h"
#include "llvm/Pass.h"
@@ -51,6 +55,120 @@ char NVPTXLowerAlloca::ID = 0;
INITIALIZE_PASS(NVPTXLowerAlloca, "nvptx-lower-alloca", "Lower Alloca", false,
false)
+static Value *getOrCreateGenericPtr(Value *LocalPtr,
+ DenseMap<Value *, Value *> &GenericPtrs) {
+ auto It = GenericPtrs.find(LocalPtr);
+ if (It != GenericPtrs.end())
+ return It->second;
+
+ auto *LocalInst = cast<Instruction>(LocalPtr);
+ auto *GenericPtr = new AddrSpaceCastInst(
+ LocalPtr, PointerType::get(LocalPtr->getContext(), ADDRESS_SPACE_GENERIC),
+ "");
+ GenericPtr->insertAfter(LocalInst->getIterator());
+ GenericPtrs[LocalPtr] = GenericPtr;
+ return GenericPtr;
+}
+
+static void updateMemIntrinsicDeclaration(MemIntrinsic *MI) {
+ SmallVector<Type *, 3> Tys;
+ if (auto *MTI = dyn_cast<MemTransferInst>(MI)) {
+ Tys.push_back(MTI->getRawDest()->getType());
+ Tys.push_back(MTI->getRawSource()->getType());
+ Tys.push_back(MTI->getLength()->getType());
+ } else {
+ auto *MSI = cast<MemSetInst>(MI);
+ Tys.push_back(MSI->getRawDest()->getType());
+ Tys.push_back(MSI->getLength()->getType());
+ }
+
+ Function *Decl =
+ Intrinsic::getOrInsertDeclaration(MI->getModule(), MI->getIntrinsicID(),
+ Tys);
+ MI->setCalledFunction(Decl);
+}
+
+static void convertPointerUsersToLocal(Value *OldPtr, Value *LocalPtr,
+ DenseMap<Value *, Value *> &GenericPtrs,
+ SmallPtrSetImpl<Value *> &Visited) {
+ if (!Visited.insert(OldPtr).second)
+ return;
+
+ for (Use &U : llvm::make_early_inc_range(OldPtr->uses())) {
+ auto *UserInst = dyn_cast<Instruction>(U.getUser());
+ if (!UserInst)
+ continue;
+
+ if (auto *LI = dyn_cast<LoadInst>(UserInst);
+ LI && LI->getPointerOperand() == OldPtr && !LI->isVolatile()) {
+ U.set(LocalPtr);
+ continue;
+ }
+
+ if (auto *SI = dyn_cast<StoreInst>(UserInst);
+ SI && SI->getPointerOperand() == OldPtr && !SI->isVolatile()) {
+ U.set(LocalPtr);
+ continue;
+ }
+
+ if (auto *GEP = dyn_cast<GetElementPtrInst>(UserInst);
+ GEP && GEP->getPointerOperand() == OldPtr) {
+ SmallVector<Value *, 4> Indices(GEP->idx_begin(), GEP->idx_end());
+ auto *NewGEP = GetElementPtrInst::Create(
+ GEP->getSourceElementType(), LocalPtr, Indices, "",
+ GEP->getIterator());
+ NewGEP->setNoWrapFlags(GEP->getNoWrapFlags());
+ NewGEP->copyMetadata(*GEP);
+ NewGEP->setDebugLoc(GEP->getDebugLoc());
+ convertPointerUsersToLocal(GEP, NewGEP, GenericPtrs, Visited);
+ if (GEP->use_empty())
+ GEP->eraseFromParent();
+ continue;
+ }
+
+ if (auto *BC = dyn_cast<BitCastInst>(UserInst);
+ BC && BC->getOperand(0) == OldPtr && BC->getType()->isPointerTy()) {
+ convertPointerUsersToLocal(BC, LocalPtr, GenericPtrs, Visited);
+ if (BC->use_empty())
+ BC->eraseFromParent();
+ continue;
+ }
+
+ if (auto *ASC = dyn_cast<AddrSpaceCastInst>(UserInst);
+ ASC && ASC->getOperand(0) == OldPtr &&
+ ASC->getDestAddressSpace() == ADDRESS_SPACE_LOCAL) {
+ ASC->replaceAllUsesWith(LocalPtr);
+ ASC->eraseFromParent();
+ continue;
+ }
+
+ if (auto *MI = dyn_cast<MemIntrinsic>(UserInst)) {
+ if (auto *MTI = dyn_cast<MemTransferInst>(MI)) {
+ bool Changed = false;
+ if (MTI->getRawDest() == OldPtr) {
+ MTI->getRawDestUse().set(LocalPtr);
+ Changed = true;
+ }
+ if (MTI->getRawSource() == OldPtr) {
+ MTI->getRawSourceUse().set(LocalPtr);
+ Changed = true;
+ }
+ if (Changed) {
+ updateMemIntrinsicDeclaration(MI);
+ continue;
+ }
+ } else if (auto *MSI = dyn_cast<MemSetInst>(MI);
+ MSI && MSI->getRawDest() == OldPtr) {
+ MSI->getRawDestUse().set(LocalPtr);
+ updateMemIntrinsicDeclaration(MI);
+ continue;
+ }
+ }
+
+ U.set(getOrCreateGenericPtr(LocalPtr, GenericPtrs));
+ }
+}
+
// =============================================================================
// Main function for this pass.
// =============================================================================
@@ -59,74 +177,87 @@ bool NVPTXLowerAlloca::runOnFunction(Function &F) {
return false;
bool Changed = false;
- for (auto &BB : F)
+ SmallVector<AllocaInst *, 8> Allocas;
+ for (auto &BB : F) {
for (auto &I : BB) {
- if (auto allocaInst = dyn_cast<AllocaInst>(&I)) {
- Changed = true;
-
- unsigned AllocAddrSpace = allocaInst->getAddressSpace();
- assert((AllocAddrSpace == ADDRESS_SPACE_GENERIC ||
- AllocAddrSpace == ADDRESS_SPACE_LOCAL) &&
- "AllocaInst can only be in Generic or Local address space for "
- "NVPTX.");
-
- Instruction *AllocaInLocalAS = allocaInst;
-
- // We need to make sure that LLVM has info that alloca needs to go to
- // ADDRESS_SPACE_LOCAL for InferAddressSpace pass.
- //
- // For allocas in ADDRESS_SPACE_GENERIC, we add addrspacecast to
- // ADDRESS_SPACE_LOCAL and back to ADDRESS_SPACE_GENERIC, so that
- // the alloca's users still use a generic pointer to operate on.
- //
- // For allocas already in ADDRESS_SPACE_LOCAL, we just need
- // addrspacecast to ADDRESS_SPACE_GENERIC.
- if (AllocAddrSpace == ADDRESS_SPACE_GENERIC) {
- auto ASCastToLocalAS = new AddrSpaceCastInst(
- allocaInst,
- PointerType::get(allocaInst->getContext(), ADDRESS_SPACE_LOCAL),
- "");
- ASCastToLocalAS->insertAfter(allocaInst->getIterator());
- AllocaInLocalAS = ASCastToLocalAS;
- }
+ if (auto *AI = dyn_cast<AllocaInst>(&I))
+ Allocas.push_back(AI);
+ }
+ }
- auto AllocaInGenericAS = new AddrSpaceCastInst(
- AllocaInLocalAS,
- PointerType::get(allocaInst->getContext(), ADDRESS_SPACE_GENERIC),
- "");
- AllocaInGenericAS->insertAfter(AllocaInLocalAS->getIterator());
-
- for (Use &AllocaUse : llvm::make_early_inc_range(allocaInst->uses())) {
- // Check Load, Store, GEP, and BitCast Uses on alloca and make them
- // use the converted generic address, in order to expose non-generic
- // addrspacecast to NVPTXInferAddressSpaces. For other types
- // of instructions this is unnecessary and may introduce redundant
- // address cast.
- auto LI = dyn_cast<LoadInst>(AllocaUse.getUser());
- if (LI && LI->getPointerOperand() == allocaInst &&
- !LI->isVolatile()) {
- LI->setOperand(LI->getPointerOperandIndex(), AllocaInGenericAS);
- continue;
- }
- auto SI = dyn_cast<StoreInst>(AllocaUse.getUser());
- if (SI && SI->getPointerOperand() == allocaInst &&
- !SI->isVolatile()) {
- SI->setOperand(SI->getPointerOperandIndex(), AllocaInGenericAS);
- continue;
- }
- auto GI = dyn_cast<GetElementPtrInst>(AllocaUse.getUser());
- if (GI && GI->getPointerOperand() == allocaInst) {
- GI->setOperand(GI->getPointerOperandIndex(), AllocaInGenericAS);
- continue;
- }
- auto BI = dyn_cast<BitCastInst>(AllocaUse.getUser());
- if (BI && BI->getOperand(0) == allocaInst) {
- BI->setOperand(0, AllocaInGenericAS);
- continue;
- }
- }
+ for (AllocaInst *allocaInst : Allocas) {
+ Changed = true;
+
+ unsigned AllocAddrSpace = allocaInst->getAddressSpace();
+ assert((AllocAddrSpace == ADDRESS_SPACE_GENERIC ||
+ AllocAddrSpace == ADDRESS_SPACE_LOCAL) &&
+ "AllocaInst can only be in Generic or Local address space for "
+ "NVPTX.");
+
+ Instruction *AllocaInLocalAS = allocaInst;
+
+ // We need to make sure that LLVM has info that alloca needs to go to
+ // ADDRESS_SPACE_LOCAL for InferAddressSpace pass.
+ //
+ // For allocas in ADDRESS_SPACE_GENERIC, create an equivalent local alloca
+ // and rewrite localizable users to use it directly. Users that still need a
+ // generic pointer use a single addrspacecast from the local alloca.
+ //
+ // For allocas already in ADDRESS_SPACE_LOCAL, we just need
+ // addrspacecast to ADDRESS_SPACE_GENERIC.
+ if (AllocAddrSpace == ADDRESS_SPACE_GENERIC) {
+ auto *AllocaInLocalAS = new AllocaInst(
+ allocaInst->getAllocatedType(), ADDRESS_SPACE_LOCAL,
+ allocaInst->getArraySize(), allocaInst->getAlign(), "",
+ allocaInst->getIterator());
+ AllocaInLocalAS->takeName(allocaInst);
+ AllocaInLocalAS->setDebugLoc(allocaInst->getDebugLoc());
+ AllocaInLocalAS->copyMetadata(*allocaInst);
+ AllocaInLocalAS->setUsedWithInAlloca(allocaInst->isUsedWithInAlloca());
+ AllocaInLocalAS->setSwiftError(allocaInst->isSwiftError());
+ DenseMap<Value *, Value *> GenericPtrs;
+ SmallPtrSet<Value *, 8> Visited;
+ convertPointerUsersToLocal(allocaInst, AllocaInLocalAS, GenericPtrs,
+ Visited);
+ assert(allocaInst->use_empty() &&
+ "All generic alloca uses should have been rewritten.");
+ allocaInst->eraseFromParent();
+ continue;
+ }
+
+ auto AllocaInGenericAS = new AddrSpaceCastInst(
+ AllocaInLocalAS,
+ PointerType::get(allocaInst->getContext(), ADDRESS_SPACE_GENERIC), "");
+ AllocaInGenericAS->insertAfter(AllocaInLocalAS->getIterator());
+
+ for (Use &AllocaUse : llvm::make_early_inc_range(allocaInst->uses())) {
+ // Check Load, Store, GEP, and BitCast Uses on alloca and make them
+ // use the converted generic address, in order to expose non-generic
+ // addrspacecast to NVPTXInferAddressSpaces. For other types
+ // of instructions this is unnecessary and may introduce redundant
+ // address cast.
+ auto LI = dyn_cast<LoadInst>(AllocaUse.getUser());
+ if (LI && LI->getPointerOperand() == allocaInst && !LI->isVolatile()) {
+ LI->setOperand(LI->getPointerOperandIndex(), AllocaInGenericAS);
+ continue;
+ }
+ auto SI = dyn_cast<StoreInst>(AllocaUse.getUser());
+ if (SI && SI->getPointerOperand() == allocaInst && !SI->isVolatile()) {
+ SI->setOperand(SI->getPointerOperandIndex(), AllocaInGenericAS);
+ continue;
+ }
+ auto GI = dyn_cast<GetElementPtrInst>(AllocaUse.getUser());
+ if (GI && GI->getPointerOperand() == allocaInst) {
+ GI->setOperand(GI->getPointerOperandIndex(), AllocaInGenericAS);
+ continue;
+ }
+ auto BI = dyn_cast<BitCastInst>(AllocaUse.getUser());
+ if (BI && BI->getOperand(0) == allocaInst) {
+ BI->setOperand(0, AllocaInGenericAS);
+ continue;
}
}
+ }
return Changed;
}
diff --git a/llvm/test/CodeGen/NVPTX/lower-alloca.ll b/llvm/test/CodeGen/NVPTX/lower-alloca.ll
index 57c1e5826c89a..d95619d5c842a 100644
--- a/llvm/test/CodeGen/NVPTX/lower-alloca.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-alloca.ll
@@ -10,12 +10,15 @@ define ptx_kernel void @kernel() {
; LABEL: @lower_alloca
; PTX-LABEL: .visible .entry kernel(
%A = alloca i32
-; CHECK: addrspacecast ptr %A to ptr addrspace(5)
-; CHECK: store i32 0, ptr addrspace(5) {{%.+}}
-; LOWERALLOCAONLY: [[V1:%.*]] = addrspacecast ptr %A to ptr addrspace(5)
-; LOWERALLOCAONLY: [[V2:%.*]] = addrspacecast ptr addrspace(5) [[V1]] to ptr
-; LOWERALLOCAONLY: store i32 0, ptr [[V2]], align 4
-; PTX: st.local.b32 [{{%rd[0-9]+}}], 0
+; CHECK: [[A:%.*]] = alloca i32, align 4, addrspace(5)
+; CHECK: [[GENERIC:%.*]] = addrspacecast ptr addrspace(5) [[A]] to ptr
+; CHECK: store i32 0, ptr addrspace(5) [[A]]
+; CHECK: call void @callee(ptr [[GENERIC]])
+; LOWERALLOCAONLY: [[A:%.*]] = alloca i32, align 4, addrspace(5)
+; LOWERALLOCAONLY: [[GENERIC:%.*]] = addrspacecast ptr addrspace(5) [[A]] to ptr
+; LOWERALLOCAONLY: store i32 0, ptr addrspace(5) [[A]], align 4
+; LOWERALLOCAONLY: call void @callee(ptr [[GENERIC]])
+; PTX: st.local.b32 [{{.*}}], 0
store i32 0, ptr %A
call void @callee(ptr %A)
ret void
diff --git a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
index 74060b7d9d9ec..6700fa3c3f4e6 100644
--- a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
@@ -140,24 +140,24 @@ define dso_local ptx_kernel void @escape_ptr(ptr nocapture noundef readnone %out
;
; PTX-LABEL: escape_ptr(
; PTX: {
-; PTX-NEXT: .local .align 4 .b8 __local_depot2[8];
+; PTX-NEXT: .local .align 8 .b8 __local_depot2[8];
; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
-; PTX-NEXT: .reg .b32 %r<3>;
-; PTX-NEXT: .reg .b64 %rd<3>;
+; PTX-NEXT: .reg .b64 %rd<7>;
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %entry
; PTX-NEXT: mov.b64 %SPL, __local_depot2;
; PTX-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-NEXT: add.u64 %rd1, %SP, 0;
-; PTX-NEXT: add.u64 %rd2, %SPL, 0;
-; PTX-NEXT: ld.param.b32 %r1, [escape_ptr_param_1+4];
-; PTX-NEXT: st.local.b32 [%rd2+4], %r1;
-; PTX-NEXT: ld.param.b32 %r2, [escape_ptr_param_1];
-; PTX-NEXT: st.local.b32 [%rd2], %r2;
+; PTX-NEXT: cvta.local.u64 %rd2, %rd1;
+; PTX-NEXT: ld.param.b32 %rd3, [escape_ptr_param_1+4];
+; PTX-NEXT: shl.b64 %rd4, %rd3, 32;
+; PTX-NEXT: ld.param.b32 %rd5, [escape_ptr_param_1];
+; PTX-NEXT: or.b64 %rd6, %rd4, %rd5;
+; PTX-NEXT: st.local.b64 [%SP], %rd6;
; PTX-NEXT: { // callseq 0, 0
; PTX-NEXT: .param .b64 param0;
-; PTX-NEXT: st.param.b64 [param0], %rd1;
+; PTX-NEXT: st.param.b64 [param0], %rd2;
; PTX-NEXT: call.uni _Z6escapePv, (param0);
; PTX-NEXT: } // callseq 0
; PTX-NEXT: ret;
@@ -188,25 +188,25 @@ define dso_local ptx_kernel void @escape_ptr_gep(ptr nocapture noundef readnone
;
; PTX-LABEL: escape_ptr_gep(
; PTX: {
-; PTX-NEXT: .local .align 4 .b8 __local_depot3[8];
+; PTX-NEXT: .local .align 8 .b8 __local_depot3[8];
; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
-; PTX-NEXT: .reg .b32 %r<3>;
-; PTX-NEXT: .reg .b64 %rd<4>;
+; PTX-NEXT: .reg .b64 %rd<8>;
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %entry
; PTX-NEXT: mov.b64 %SPL, __local_depot3;
; PTX-NEXT: cvta.local.u64 %SP, %SPL;
-; PTX-NEXT: add.u64 %rd1, %SP, 0;
-; PTX-NEXT: add.u64 %rd2, %SPL, 0;
-; PTX-NEXT: ld.param.b32 %r1, [escape_ptr_gep_param_1+4];
-; PTX-NEXT: st.local.b32 [%rd2+4], %r1;
-; PTX-NEXT: ld.param.b32 %r2, [escape_ptr_gep_param_1];
-; PTX-NEXT: st.local.b32 [%rd2], %r2;
-; PTX-NEXT: add.s64 %rd3, %rd1, 4;
+; PTX-NEXT: ld.param.b32 %rd1, [escape_ptr_gep_param_1+4];
+; PTX-NEXT: shl.b64 %rd2, %rd1, 32;
+; PTX-NEXT: ld.param.b32 %rd3, [escape_ptr_gep_param_1];
+; PTX-NEXT: or.b64 %rd4, %rd2, %rd3;
+; PTX-NEXT: st.local.b64 [%SP], %rd4;
+; PTX-NEXT: add.u64 %rd5, %SP, 0;
+; PTX-NEXT: or.b64 %rd6, %rd5, 4;
+; PTX-NEXT: cvta.local.u64 %rd7, %rd6;
; PTX-NEXT: { // callseq 1, 0
; PTX-NEXT: .param .b64 param0;
-; PTX-NEXT: st.param.b64 [param0], %rd3;
+; PTX-NEXT: st.param.b64 [param0], %rd7;
; PTX-NEXT: call.uni _Z6escapePv, (param0);
; PTX-NEXT: } // callseq 1
; PTX-NEXT: ret;
@@ -236,11 +236,10 @@ define dso_local ptx_kernel void @escape_ptr_store(ptr nocapture noundef writeon
;
; PTX-LABEL: escape_ptr_store(
; PTX: {
-; PTX-NEXT: .local .align 4 .b8 __local_depot4[8];
+; PTX-NEXT: .local .align 8 .b8 __local_depot4[8];
; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
-; PTX-NEXT: .reg .b32 %r<3>;
-; PTX-NEXT: .reg .b64 %rd<5>;
+; PTX-NEXT: .reg .b64 %rd<9>;
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %entry
; PTX-NEXT: mov.b64 %SPL, __local_depot4;
@@ -248,12 +247,13 @@ define dso_local ptx_kernel void @escape_ptr_store(ptr nocapture noundef writeon
; PTX-NEXT: ld.param.b64 %rd1, [escape_ptr_store_param_0];
; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;
; PTX-NEXT: add.u64 %rd3, %SP, 0;
-; PTX-NEXT: add.u64 %rd4, %SPL, 0;
-; PTX-NEXT: ld.param.b32 %r1, [escape_ptr_store_param_1+4];
-; PTX-NEXT: st.local.b32 [%rd4+4], %r1;
-; PTX-NEXT: ld.param.b32 %r2, [escape_ptr_store_param_1];
-; PTX-NEXT: st.local.b32 [%rd4], %r2;
-; PTX-NEXT: st.global.b64 [%rd2], %rd3;
+; PTX-NEXT: cvta.local.u64 %rd4, %rd3;
+; PTX-NEXT: ld.param.b32 %rd5, [escape_ptr_store_param_1+4];
+; PTX-NEXT: shl.b64 %rd6, %rd5, 32;
+; PTX-NEXT: ld.param.b32 %rd7, [escape_ptr_store_param_1];
+; PTX-NEXT: or.b64 %rd8, %rd6, %rd7;
+; PTX-NEXT: st.local.b64 [%SP], %rd8;
+; PTX-NEXT: st.global.b64 [%rd2], %rd4;
; PTX-NEXT: ret;
entry:
store ptr %s, ptr %out, align 8
@@ -282,25 +282,25 @@ define dso_local ptx_kernel void @escape_ptr_gep_store(ptr nocapture noundef wri
;
; PTX-LABEL: escape_ptr_gep_store(
; PTX: {
-; PTX-NEXT: .local .align 4 .b8 __local_depot5[8];
+; PTX-NEXT: .local .align 8 .b8 __local_depot5[8];
; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
-; PTX-NEXT: .reg .b32 %r<3>;
-; PTX-NEXT: .reg .b64 %rd<6>;
+; PTX-NEXT: .reg .b64 %rd<10>;
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %entry
; PTX-NEXT: mov.b64 %SPL, __local_depot5;
; PTX-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-NEXT: ld.param.b64 %rd1, [escape_ptr_gep_store_param_0];
; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;
-; PTX-NEXT: add.u64 %rd3, %SP, 0;
-; PTX-NEXT: add.u64 %rd4, %SPL, 0;
-; PTX-NEXT: ld.param.b32 %r1, [escape_ptr_gep_store_param_1+4];
-; PTX-NEXT: st.local.b32 [%rd4+4], %r1;
-; PTX-NEXT: ld.param.b32 %r2, [escape_ptr_gep_store_param_1];
-; PTX-NEXT: st.local.b32 [%rd4], %r2;
-; PTX-NEXT: add.s64 %rd5, %rd3, 4;
-; PTX-NEXT: st.global.b64 [%rd2], %rd5;
+; PTX-NEXT: ld.param.b32 %rd3, [escape_ptr_gep_store_param_1+4];
+; PTX-NEXT: shl.b64 %rd4, %rd3, 32;
+; PTX-NEXT: ld.param.b32 %rd5, [escape_ptr_gep_store_param_1];
+; PTX-NEXT: or.b64 %rd6, %rd4, %rd5;
+; PTX-NEXT: st.local.b64 [%SP], %rd6;
+; PTX-NEXT: add.u64 %rd7, %SP, 0;
+; PTX-NEXT: or.b64 %rd8, %rd7, 4;
+; PTX-NEXT: cvta.local.u64 %rd9, %rd8;
+; PTX-NEXT: st.global.b64 [%rd2], %rd9;
; PTX-NEXT: ret;
entry:
%b = getelementptr inbounds nuw i8, ptr %s, i64 4
@@ -330,11 +330,10 @@ define dso_local ptx_kernel void @escape_ptrtoint(ptr nocapture noundef writeonl
;
; PTX-LABEL: escape_ptrtoint(
; PTX: {
-; PTX-NEXT: .local .align 4 .b8 __local_depot6[8];
+; PTX-NEXT: .local .align 8 .b8 __local_depot6[8];
; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
-; PTX-NEXT: .reg .b32 %r<3>;
-; PTX-NEXT: .reg .b64 %rd<5>;
+; PTX-NEXT: .reg .b64 %rd<9>;
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %entry
; PTX-NEXT: mov.b64 %SPL, __local_depot6;
@@ -342,12 +341,13 @@ define dso_local ptx_kernel void @escape_ptrtoint(ptr nocapture noundef writeonl
; PTX-NEXT: ld.param.b64 %rd1, [escape_ptrtoint_param_0];
; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;
; PTX-NEXT: add.u64 %rd3, %SP, 0;
-; PTX-NEXT: add.u64 %rd4, %SPL, 0;
-; PTX-NEXT: ld.param.b32 %r1, [escape_ptrtoint_param_1+4];
-; PTX-NEXT: st.local.b32 [%rd4+4], %r1;
-; PTX-NEXT: ld.param.b32 %r2, [escape_ptrtoint_param_1];
-; PTX-NEXT: st.local.b32 [%rd4], %r2;
-; PTX-NEXT: st.global.b64 [%rd2], %rd3;
+; PTX-NEXT: cvta.local.u64 %rd4, %rd3;
+; PTX-NEXT: ld.param.b32 %rd5, [escape_ptrtoint_param_1+4];
+; PTX-NEXT: shl.b64 %rd6, %rd5, 32;
+; PTX-NEXT: ld.param.b32 %rd7, [escape_ptrtoint_param_1];
+; PTX-NEXT: or.b64 %rd8, %rd6, %rd7;
+; PTX-NEXT: st.local.b64 [%SP], %rd8;
+; PTX-NEXT: st.global.b64 [%rd2], %rd4;
; PTX-NEXT: ret;
entry:
%i = ptrtoint ptr %s to i64
@@ -497,7 +497,7 @@ define dso_local ptx_kernel void @memcpy_to_param(ptr nocapture noundef readonly
; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
; PTX-NEXT: .reg .b32 %r<23>;
-; PTX-NEXT: .reg .b64 %rd<4>;
+; PTX-NEXT: .reg .b64 %rd<3>;
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %entry
; PTX-NEXT: mov.b64 %SPL, __local_depot9;
@@ -506,7 +506,7 @@ define dso_local ptx_kernel void @memcpy_to_param(ptr nocapture noundef readonly
; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;
; PTX-NEXT: ld.param.b32 %r1, [memcpy_to_param_param_1+4];
; PTX-NEXT: ld.param.b32 %r2, [memcpy_to_param_param_1];
-; PTX-NEXT: st.v2.b32 [%SP], {%r2, %r1};
+; PTX-NEXT: st.local.v2.b32 [%SP], {%r2, %r1};
; PTX-NEXT: ld.volatile.global.b8 %r3, [%rd2+4];
; PTX-NEXT: ld.volatile.global.b8 %r4, [%rd2+5];
; PTX-NEXT: shl.b32 %r5, %r4, 8;
@@ -527,8 +527,7 @@ define dso_local ptx_kernel void @memcpy_to_param(ptr nocapture noundef readonly
; PTX-NEXT: shl.b32 %r20, %r19, 24;
; PTX-NEXT: or.b32 %r21, %r20, %r18;
; PTX-NEXT: or.b32 %r22, %r21, %r16;
-; PTX-NEXT: add.u64 %rd3, %SPL, 0;
-; PTX-NEXT: st.local.v2.b32 [%rd3], {%r22, %r12};
+; PTX-NEXT: st.local.v2.b32 [%SP], {%r22, %r12};
; PTX-NEXT: ret;
entry:
tail call void @llvm.memcpy.p0.p0.i64(ptr %s, ptr %in, i64 16, i1 true)
@@ -668,12 +667,12 @@ define ptx_kernel void @test_select_write(ptr byval(i32) align 4 %input1, ptr by
; PTX-NEXT: and.b16 %rs2, %rs1, 1;
; PTX-NEXT: setp.ne.b16 %p1, %rs2, 0;
; PTX-NEXT: ld.param.b32 %r1, [test_select_write_param_1];
-; PTX-NEXT: st.b32 [%SP], %r1;
+; PTX-NEXT: st.local.b32 [%SP], %r1;
; PTX-NEXT: ld.param.b32 %r2, [test_select_write_param_0];
-; PTX-NEXT: st.b32 [%SP+4], %r2;
-; PTX-NEXT: add.u64 %rd1, %SPL, 4;
-; PTX-NEXT: add.u64 %rd2, %SPL, 0;
-; PTX-NEXT: selp.b64 %rd3, %rd1, %rd2, %p1;
+; PTX-NEXT: st.local.b32 [%SP+4], %r2;
+; PTX-NEXT: add.u64 %rd1, %SP, 0;
+; PTX-NEXT: add.u64 %rd2, %SP, 4;
+; PTX-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1;
; PTX-NEXT: st.local.b32 [%rd3], 1;
; PTX-NEXT: ret;
bb:
@@ -858,13 +857,13 @@ define ptx_kernel void @test_phi_write(ptr byval(%struct.S) align 4 %input1, ptr
; PTX-NEXT: and.b16 %rs2, %rs1, 1;
; PTX-NEXT: setp.ne.b16 %p1, %rs2, 0;
; PTX-NEXT: ld.param.b32 %r1, [test_phi_write_param_1+4];
-; PTX-NEXT: st.b32 [%SP], %r1;
-; PTX-NEXT: add.u64 %rd1, %SPL, 4;
+; PTX-NEXT: st.local.b32 [%SP], %r1;
; PTX-NEXT: ld.param.b32 %r2, [test_phi_write_param_0];
-; PTX-NEXT: st.b32 [%SP+4], %r2;
+; PTX-NEXT: st.local.b32 [%SP+4], %r2;
+; PTX-NEXT: add.u64 %rd1, %SP, 4;
; PTX-NEXT: @%p1 bra $L__BB14_2;
; PTX-NEXT: // %bb.1: // %second
-; PTX-NEXT: add.u64 %rd1, %SPL, 0;
+; PTX-NEXT: add.u64 %rd1, %SP, 0;
; PTX-NEXT: $L__BB14_2: // %merge
; PTX-NEXT: st.local.b32 [%rd1], 1;
; PTX-NEXT: ret;
@@ -905,17 +904,17 @@ define ptx_kernel void @test_forward_byval_arg(ptr byval(i32) align 4 %input) {
; PTX-NEXT: .local .align 4 .b8 __local_depot15[4];
; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
-; PTX-NEXT: .reg .b32 %r<2>;
-; PTX-NEXT: .reg .b64 %rd<2>;
+; PTX-NEXT: .reg .b32 %r<3>;
; PTX-EMPTY:
; PTX-NEXT: // %bb.0:
; PTX-NEXT: mov.b64 %SPL, __local_depot15;
-; PTX-NEXT: add.u64 %rd1, %SPL, 0;
+; PTX-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-NEXT: ld.param.b32 %r1, [test_forward_byval_arg_param_0];
-; PTX-NEXT: st.local.b32 [%rd1], %r1;
+; PTX-NEXT: st.local.b32 [%SP], %r1;
; PTX-NEXT: { // callseq 2, 0
; PTX-NEXT: .param .align 4 .b8 param0[4];
-; PTX-NEXT: st.param.b32 [param0], %r1;
+; PTX-NEXT: ld.b32 %r2, [%SP];
+; PTX-NEXT: st.param.b32 [param0], %r2;
; PTX-NEXT: call.uni device_func, (param0);
; PTX-NEXT: } // callseq 2
; PTX-NEXT: ret;
diff --git a/llvm/test/CodeGen/NVPTX/memcpy-alloca-align.ll b/llvm/test/CodeGen/NVPTX/memcpy-alloca-align.ll
new file mode 100644
index 0000000000000..601ce9477bd69
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/memcpy-alloca-align.ll
@@ -0,0 +1,132 @@
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_50 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_50 | %ptxas-verify %}
+
+; The expansion of small memcpy/memmove/memset can raise the alignment of
+; under-aligned stack objects. Verify this also works on NVPTX, which lowers
+; allocas to the local address space via addrspacecast instructions that hide
+; the frame index, both when the pointer is materialized in the same basic
+; block as the memory operation and when it is defined in a different one.
+
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
+target triple = "nvptx64-nvidia-cuda"
+
+declare void @use(ptr)
+
+define void @memcpy_alloca() {
+; CHECK-LABEL: memcpy_alloca(
+; CHECK: .local .align 4 .b8 __local_depot{{[0-9]+}}[16];
+; CHECK: st.local.b8
+; CHECK: st.local.b16
+; CHECK: st.local.b32
+; CHECK-NOT: st.local.b8
+; CHECK: ret;
+ %a = alloca [7 x i8], align 1
+ %b = alloca [7 x i8], align 1
+ call void @use(ptr %a)
+ call void @use(ptr %b)
+ call void @llvm.memcpy.p0.p0.i64(ptr align 1 %a, ptr align 1 %b, i64 7, i1 false)
+ call void @use(ptr %a)
+ ret void
+}
+
+define void @memcpy_alloca_cross_bb(i1 %c) {
+; CHECK-LABEL: memcpy_alloca_cross_bb(
+; CHECK: .local .align 4 .b8 __local_depot{{[0-9]+}}[16];
+; CHECK: st.local.b8
+; CHECK: st.local.b16
+; CHECK: st.local.b32
+; CHECK-NOT: st.local.b8
+; CHECK: ret;
+entry:
+ %a = alloca [7 x i8], align 1
+ %b = alloca [7 x i8], align 1
+ call void @use(ptr %a)
+ call void @use(ptr %b)
+ br i1 %c, label %copy, label %exit
+
+copy:
+ call void @llvm.memcpy.p0.p0.i64(ptr align 1 %a, ptr align 1 %b, i64 7, i1 false)
+ call void @use(ptr %a)
+ br label %exit
+
+exit:
+ ret void
+}
+
+; The source object's alignment was raised when expanding the first memcpy;
+; the loads of the second copy can only learn about it through the frame index.
+define void @memcpy_alloca_raised_src(i1 %c, ptr %p) {
+; CHECK-LABEL: memcpy_alloca_raised_src(
+; CHECK: .local .align 4 .b8 __local_depot{{[0-9]+}}[16];
+; CHECK: ld.local.b8
+; CHECK: ld.local.b16
+; CHECK: ld.local.b32
+; CHECK-NOT: ld.local.b8
+; CHECK: ret;
+entry:
+ %a = alloca [7 x i8], align 1
+ %b = alloca [7 x i8], align 1
+ call void @use(ptr %a)
+ call void @use(ptr %b)
+ br i1 %c, label %copy1, label %exit
+
+copy1:
+ call void @llvm.memcpy.p0.p0.i64(ptr align 1 %b, ptr align 1 %p, i64 7, i1 false)
+ br label %copy2
+
+copy2:
+ call void @llvm.memcpy.p0.p0.i64(ptr align 1 %a, ptr align 1 %b, i64 7, i1 false)
+ call void @use(ptr %a)
+ br label %exit
+
+exit:
+ ret void
+}
+
+define void @memmove_alloca() {
+; CHECK-LABEL: memmove_alloca(
+; CHECK: .local .align 4 .b8 __local_depot{{[0-9]+}}[16];
+; CHECK: st.local.b8
+; CHECK: st.local.b16
+; CHECK: st.local.b32
+; CHECK-NOT: st.local.b8
+; CHECK: ret;
+ %a = alloca [7 x i8], align 1
+ %b = alloca [7 x i8], align 1
+ call void @use(ptr %a)
+ call void @use(ptr %b)
+ call void @llvm.memmove.p0.p0.i64(ptr align 1 %a, ptr align 1 %b, i64 7, i1 false)
+ call void @use(ptr %a)
+ ret void
+}
+
+define void @memset_alloca() {
+; CHECK-LABEL: memset_alloca(
+; CHECK: .local .align 4 .b8 __local_depot{{[0-9]+}}[8];
+; CHECK: st.local.b8
+; CHECK: st.local.b16
+; CHECK: st.local.b32
+; CHECK-NOT: st.local.b8
+; CHECK: ret;
+ %a = alloca [7 x i8], align 1
+ call void @use(ptr %a)
+ call void @llvm.memset.p0.i64(ptr align 1 %a, i8 0, i64 7, i1 false)
+ call void @use(ptr %a)
+ ret void
+}
+
+; InferPtrAlign also looks through the addrspacecast, letting DAGCombiner
+; refine the alignment of under-aligned plain loads and stores.
+define i32 @underaligned_load_store(i32 %v) {
+; CHECK-LABEL: underaligned_load_store(
+; CHECK: ld.local.b32
+; CHECK: st.local.b32
+; CHECK-NOT: .local.b8
+; CHECK: ret;
+ %a = alloca i32, align 4
+ call void @use(ptr %a)
+ %l = load i32, ptr %a, align 1
+ store i32 %v, ptr %a, align 1
+ call void @use(ptr %a)
+ ret i32 %l
+}
diff --git a/llvm/test/CodeGen/NVPTX/variadics-backend.ll b/llvm/test/CodeGen/NVPTX/variadics-backend.ll
index a077d29d092f6..181d38093b0ba 100644
--- a/llvm/test/CodeGen/NVPTX/variadics-backend.ll
+++ b/llvm/test/CodeGen/NVPTX/variadics-backend.ll
@@ -109,17 +109,17 @@ define dso_local i32 @foo() {
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot1;
-; CHECK-PTX-NEXT: add.u64 %rd1, %SPL, 0;
-; CHECK-PTX-NEXT: st.local.b32 [%rd1], 1;
-; CHECK-PTX-NEXT: st.local.b32 [%rd1+4], 1;
-; CHECK-PTX-NEXT: st.local.b32 [%rd1+8], 1;
-; CHECK-PTX-NEXT: st.local.b64 [%rd1+16], 1;
-; CHECK-PTX-NEXT: st.local.b64 [%rd1+24], 4607182418800017408;
-; CHECK-PTX-NEXT: st.local.b64 [%rd1+32], 4607182418800017408;
+; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;
+; CHECK-PTX-NEXT: st.local.b64 [%SP], 4294967297;
+; CHECK-PTX-NEXT: st.local.b32 [%SP+8], 1;
+; CHECK-PTX-NEXT: st.local.b64 [%SP+16], 1;
+; CHECK-PTX-NEXT: st.local.b64 [%SP+24], 4607182418800017408;
+; CHECK-PTX-NEXT: st.local.b64 [%SP+32], 4607182418800017408;
; CHECK-PTX-NEXT: { // callseq 0, 0
; CHECK-PTX-NEXT: .param .b32 param0;
; CHECK-PTX-NEXT: .param .b64 param1;
; CHECK-PTX-NEXT: .param .b32 retval0;
+; CHECK-PTX-NEXT: add.u64 %rd1, %SP, 0;
; CHECK-PTX-NEXT: st.param.b64 [param1], %rd1;
; CHECK-PTX-NEXT: st.param.b32 [param0], 1;
; CHECK-PTX-NEXT: call.uni (retval0), variadics1, (param0, param1);
@@ -138,34 +138,35 @@ entry:
define dso_local i32 @variadics2(i32 noundef %first, ...) {
; CHECK-PTX-LABEL: variadics2(
; CHECK-PTX: {
-; CHECK-PTX-NEXT: .local .align 1 .b8 __local_depot2[3];
+; CHECK-PTX-NEXT: .local .align 2 .b8 __local_depot2[4];
; CHECK-PTX-NEXT: .reg .b64 %SP;
; CHECK-PTX-NEXT: .reg .b64 %SPL;
-; CHECK-PTX-NEXT: .reg .b16 %rs<4>;
+; CHECK-PTX-NEXT: .reg .b16 %rs<6>;
; CHECK-PTX-NEXT: .reg .b32 %r<6>;
-; CHECK-PTX-NEXT: .reg .b64 %rd<8>;
+; CHECK-PTX-NEXT: .reg .b64 %rd<7>;
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot2;
+; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;
; CHECK-PTX-NEXT: ld.param.b32 %r1, [variadics2_param_0];
; CHECK-PTX-NEXT: ld.param.b64 %rd1, [variadics2_param_1];
-; CHECK-PTX-NEXT: add.u64 %rd2, %SPL, 0;
-; CHECK-PTX-NEXT: add.s64 %rd3, %rd1, 7;
-; CHECK-PTX-NEXT: and.b64 %rd4, %rd3, -8;
-; CHECK-PTX-NEXT: ld.local.b32 %r2, [%rd4];
-; CHECK-PTX-NEXT: ld.local.s8 %r3, [%rd4+4];
-; CHECK-PTX-NEXT: ld.local.b8 %rs1, [%rd4+7];
-; CHECK-PTX-NEXT: st.local.b8 [%rd2+2], %rs1;
-; CHECK-PTX-NEXT: ld.local.b8 %rs2, [%rd4+6];
-; CHECK-PTX-NEXT: st.local.b8 [%rd2+1], %rs2;
-; CHECK-PTX-NEXT: ld.local.b8 %rs3, [%rd4+5];
-; CHECK-PTX-NEXT: st.local.b8 [%rd2], %rs3;
-; CHECK-PTX-NEXT: ld.local.b64 %rd5, [%rd4+8];
+; CHECK-PTX-NEXT: add.s64 %rd2, %rd1, 7;
+; CHECK-PTX-NEXT: and.b64 %rd3, %rd2, -8;
+; CHECK-PTX-NEXT: ld.local.b32 %r2, [%rd3];
+; CHECK-PTX-NEXT: ld.local.s8 %r3, [%rd3+4];
+; CHECK-PTX-NEXT: ld.local.b8 %rs1, [%rd3+7];
+; CHECK-PTX-NEXT: st.local.b8 [%SP+2], %rs1;
+; CHECK-PTX-NEXT: ld.local.b8 %rs2, [%rd3+5];
+; CHECK-PTX-NEXT: ld.local.b8 %rs3, [%rd3+6];
+; CHECK-PTX-NEXT: shl.b16 %rs4, %rs3, 8;
+; CHECK-PTX-NEXT: or.b16 %rs5, %rs4, %rs2;
+; CHECK-PTX-NEXT: st.local.b16 [%SP], %rs5;
+; CHECK-PTX-NEXT: ld.local.b64 %rd4, [%rd3+8];
; CHECK-PTX-NEXT: add.s32 %r4, %r1, %r2;
; CHECK-PTX-NEXT: add.s32 %r5, %r4, %r3;
-; CHECK-PTX-NEXT: cvt.u64.u32 %rd6, %r5;
-; CHECK-PTX-NEXT: add.s64 %rd7, %rd6, %rd5;
-; CHECK-PTX-NEXT: st.param.b32 [func_retval0], %rd7;
+; CHECK-PTX-NEXT: cvt.u64.u32 %rd5, %r5;
+; CHECK-PTX-NEXT: add.s64 %rd6, %rd5, %rd4;
+; CHECK-PTX-NEXT: st.param.b32 [func_retval0], %rd6;
; CHECK-PTX-NEXT: ret;
entry:
%vlist = alloca ptr, align 8
@@ -201,28 +202,29 @@ define dso_local i32 @bar() {
; CHECK-PTX-NEXT: .local .align 8 .b8 __local_depot3[24];
; CHECK-PTX-NEXT: .reg .b64 %SP;
; CHECK-PTX-NEXT: .reg .b64 %SPL;
-; CHECK-PTX-NEXT: .reg .b16 %rs<4>;
+; CHECK-PTX-NEXT: .reg .b16 %rs<6>;
; CHECK-PTX-NEXT: .reg .b32 %r<2>;
-; CHECK-PTX-NEXT: .reg .b64 %rd<3>;
+; CHECK-PTX-NEXT: .reg .b64 %rd<2>;
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot3;
-; CHECK-PTX-NEXT: add.u64 %rd1, %SPL, 0;
-; CHECK-PTX-NEXT: add.u64 %rd2, %SPL, 8;
+; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;
; CHECK-PTX-NEXT: ld.global.nc.b8 %rs1, [__const_$_bar_$_s1+7];
-; CHECK-PTX-NEXT: st.local.b8 [%rd1+2], %rs1;
+; CHECK-PTX-NEXT: st.local.b8 [%SP+2], %rs1;
; CHECK-PTX-NEXT: ld.global.nc.b8 %rs2, [__const_$_bar_$_s1+6];
-; CHECK-PTX-NEXT: st.local.b8 [%rd1+1], %rs2;
-; CHECK-PTX-NEXT: ld.global.nc.b8 %rs3, [__const_$_bar_$_s1+5];
-; CHECK-PTX-NEXT: st.local.b8 [%rd1], %rs3;
-; CHECK-PTX-NEXT: st.local.b32 [%rd2], 1;
-; CHECK-PTX-NEXT: st.local.b8 [%rd2+4], 1;
-; CHECK-PTX-NEXT: st.local.b64 [%rd2+8], 1;
+; CHECK-PTX-NEXT: shl.b16 %rs3, %rs2, 8;
+; CHECK-PTX-NEXT: ld.global.nc.b8 %rs4, [__const_$_bar_$_s1+5];
+; CHECK-PTX-NEXT: or.b16 %rs5, %rs3, %rs4;
+; CHECK-PTX-NEXT: st.local.b16 [%SP], %rs5;
+; CHECK-PTX-NEXT: st.local.b32 [%SP+8], 1;
+; CHECK-PTX-NEXT: st.local.b8 [%SP+12], 1;
+; CHECK-PTX-NEXT: st.local.b64 [%SP+16], 1;
; CHECK-PTX-NEXT: { // callseq 1, 0
; CHECK-PTX-NEXT: .param .b32 param0;
; CHECK-PTX-NEXT: .param .b64 param1;
; CHECK-PTX-NEXT: .param .b32 retval0;
-; CHECK-PTX-NEXT: st.param.b64 [param1], %rd2;
+; CHECK-PTX-NEXT: add.u64 %rd1, %SP, 8;
+; CHECK-PTX-NEXT: st.param.b64 [param1], %rd1;
; CHECK-PTX-NEXT: st.param.b32 [param0], 1;
; CHECK-PTX-NEXT: call.uni (retval0), variadics2, (param0, param1);
; CHECK-PTX-NEXT: ld.param.b32 %r1, [retval0];
@@ -286,12 +288,13 @@ define dso_local i32 @baz() {
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot5;
-; CHECK-PTX-NEXT: add.u64 %rd1, %SPL, 0;
-; CHECK-PTX-NEXT: st.local.v4.b32 [%rd1], {1, 1, 1, 1};
+; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;
+; CHECK-PTX-NEXT: st.local.v4.b32 [%SP], {1, 1, 1, 1};
; CHECK-PTX-NEXT: { // callseq 2, 0
; CHECK-PTX-NEXT: .param .b32 param0;
; CHECK-PTX-NEXT: .param .b64 param1;
; CHECK-PTX-NEXT: .param .b32 retval0;
+; CHECK-PTX-NEXT: add.u64 %rd1, %SP, 0;
; CHECK-PTX-NEXT: st.param.b64 [param1], %rd1;
; CHECK-PTX-NEXT: st.param.b32 [param0], 1;
; CHECK-PTX-NEXT: call.uni (retval0), variadics3, (param0, param1);
@@ -346,26 +349,26 @@ define dso_local void @qux() {
; CHECK-PTX-NEXT: .local .align 8 .b8 __local_depot7[24];
; CHECK-PTX-NEXT: .reg .b64 %SP;
; CHECK-PTX-NEXT: .reg .b64 %SPL;
-; CHECK-PTX-NEXT: .reg .b64 %rd<7>;
+; CHECK-PTX-NEXT: .reg .b64 %rd<6>;
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot7;
-; CHECK-PTX-NEXT: add.u64 %rd1, %SPL, 0;
-; CHECK-PTX-NEXT: add.u64 %rd2, %SPL, 16;
-; CHECK-PTX-NEXT: ld.global.nc.b64 %rd3, [__const_$_qux_$_s+8];
-; CHECK-PTX-NEXT: st.local.b64 [%rd1+8], %rd3;
-; CHECK-PTX-NEXT: ld.global.nc.b64 %rd4, [__const_$_qux_$_s];
-; CHECK-PTX-NEXT: st.local.b64 [%rd1], %rd4;
-; CHECK-PTX-NEXT: st.local.b64 [%rd2], 1;
+; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;
+; CHECK-PTX-NEXT: ld.global.nc.b64 %rd1, [__const_$_qux_$_s+8];
+; CHECK-PTX-NEXT: st.local.b64 [%SP+8], %rd1;
+; CHECK-PTX-NEXT: ld.global.nc.b64 %rd2, [__const_$_qux_$_s];
+; CHECK-PTX-NEXT: st.local.b64 [%SP], %rd2;
+; CHECK-PTX-NEXT: st.local.b64 [%SP+16], 1;
; CHECK-PTX-NEXT: { // callseq 3, 0
; CHECK-PTX-NEXT: .param .align 8 .b8 param0[16];
; CHECK-PTX-NEXT: .param .b64 param1;
; CHECK-PTX-NEXT: .param .b32 retval0;
-; CHECK-PTX-NEXT: st.param.b64 [param1], %rd2;
-; CHECK-PTX-NEXT: ld.local.b64 %rd5, [%rd1+8];
-; CHECK-PTX-NEXT: st.param.b64 [param0+8], %rd5;
-; CHECK-PTX-NEXT: ld.local.b64 %rd6, [%rd1];
-; CHECK-PTX-NEXT: st.param.b64 [param0], %rd6;
+; CHECK-PTX-NEXT: add.u64 %rd3, %SP, 16;
+; CHECK-PTX-NEXT: st.param.b64 [param1], %rd3;
+; CHECK-PTX-NEXT: ld.b64 %rd4, [%SP+8];
+; CHECK-PTX-NEXT: st.param.b64 [param0+8], %rd4;
+; CHECK-PTX-NEXT: ld.b64 %rd5, [%SP];
+; CHECK-PTX-NEXT: st.param.b64 [param0], %rd5;
; CHECK-PTX-NEXT: call.uni (retval0), variadics4, (param0, param1);
; CHECK-PTX-NEXT: } // callseq 3
; CHECK-PTX-NEXT: ret;
>From 9e28cc7671770d17fbdef736c5b72ce4543ad90a Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Wed, 17 Jun 2026 14:03:25 +0200
Subject: [PATCH 02/15] Fix stack lowering for local allocas
Resolve addrspace(5) frame indices to the local frame pointer %SPL, so
.local accesses use a local address and escaping casts fold to a single
cvta.local. Run NVPTXLowerAlloca unconditionally (now a correctness
requirement, not an optimization), drop the redundant cast in
LowerDYNAMIC_STACKALLOC, and preserve lifetime markers and debug
locations across the rewrite.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply at anthropic.com>
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 12 ++-
llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp | 26 ++++-
llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp | 14 ++-
llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp | 4 +
.../CodeGen/NVPTX/call-with-alloca-buffer.ll | 5 +-
.../CodeGen/NVPTX/convert-call-to-indirect.ll | 10 +-
.../NVPTX/dynamic-stackalloc-regression.ll | 14 +--
llvm/test/CodeGen/NVPTX/indirect_byval.ll | 22 ++---
llvm/test/CodeGen/NVPTX/local-stack-frame.ll | 73 +++++++-------
llvm/test/CodeGen/NVPTX/lower-alloca.ll | 4 +-
llvm/test/CodeGen/NVPTX/lower-byval-args.ll | 53 +++++-----
llvm/test/CodeGen/NVPTX/vaargs.ll | 96 +++++++++----------
llvm/test/CodeGen/NVPTX/variadics-backend.ll | 49 +++++-----
.../DebugInfo/NVPTX/dbg-declare-alloca.ll | 5 +-
14 files changed, 199 insertions(+), 188 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 8a622396e7f6d..cc7d52b38f5b3 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1714,10 +1714,14 @@ SDValue NVPTXTargetLowering::LowerDYNAMIC_STACKALLOC(SDValue Op,
{Chain, DAG.getZExtOrTrunc(Size, DL, LocalVT),
DAG.getTargetConstant(Align, DL, MVT::i32)});
- SDValue ASC = DAG.getAddrSpaceCast(
- DL, Op.getValueType(), Alloc, ADDRESS_SPACE_LOCAL, ADDRESS_SPACE_GENERIC);
-
- return DAG.getMergeValues({ASC, SDValue(Alloc.getNode(), 1)}, DL);
+ // Keep the result in the alloca's address space: return the local pointer as
+ // is when that is requested (escapes are cast by the caller), else cast it.
+ SDValue Ptr = Alloc;
+ if (Op.getValueType() != LocalVT)
+ Ptr = DAG.getAddrSpaceCast(DL, Op.getValueType(), Alloc,
+ ADDRESS_SPACE_LOCAL, ADDRESS_SPACE_GENERIC);
+
+ return DAG.getMergeValues({Ptr, SDValue(Alloc.getNode(), 1)}, DL);
}
SDValue NVPTXTargetLowering::LowerSTACKRESTORE(SDValue Op,
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
index 51d8b1ab3b9b1..8483d216489d6 100644
--- a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
@@ -28,6 +28,7 @@
#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/ADT/SmallVector.h"
+#include "llvm/IR/DebugInfo.h"
#include "llvm/IR/Function.h"
#include "llvm/IR/IntrinsicInst.h"
#include "llvm/IR/Intrinsics.h"
@@ -94,6 +95,13 @@ static void convertPointerUsersToLocal(Value *OldPtr, Value *LocalPtr,
if (!Visited.insert(OldPtr).second)
return;
+ // Debug records aren't on the use-list visited below, so retarget them here;
+ // otherwise the variable's location is lost when the old alloca is erased.
+ SmallVector<DbgVariableRecord *, 2> DbgUsers;
+ findDbgUsers(OldPtr, DbgUsers);
+ for (DbgVariableRecord *DVR : DbgUsers)
+ DVR->replaceVariableLocationOp(OldPtr, LocalPtr);
+
for (Use &U : llvm::make_early_inc_range(OldPtr->uses())) {
auto *UserInst = dyn_cast<Instruction>(U.getUser());
if (!UserInst)
@@ -142,6 +150,19 @@ static void convertPointerUsersToLocal(Value *OldPtr, Value *LocalPtr,
continue;
}
+ if (auto *II = dyn_cast<IntrinsicInst>(UserInst);
+ II && (II->getIntrinsicID() == Intrinsic::lifetime_start ||
+ II->getIntrinsicID() == Intrinsic::lifetime_end) &&
+ isa<AllocaInst>(LocalPtr)) {
+ // Lifetime markers must reference an alloca directly, so retarget them
+ // to the local alloca and update the overloaded declaration.
+ U.set(LocalPtr);
+ Function *Decl = Intrinsic::getOrInsertDeclaration(
+ II->getModule(), II->getIntrinsicID(), {LocalPtr->getType()});
+ II->setCalledFunction(Decl);
+ continue;
+ }
+
if (auto *MI = dyn_cast<MemIntrinsic>(UserInst)) {
if (auto *MTI = dyn_cast<MemTransferInst>(MI)) {
bool Changed = false;
@@ -173,9 +194,8 @@ static void convertPointerUsersToLocal(Value *OldPtr, Value *LocalPtr,
// Main function for this pass.
// =============================================================================
bool NVPTXLowerAlloca::runOnFunction(Function &F) {
- if (skipFunction(F))
- return false;
-
+ // Mandatory lowering: later stack lowering relies on local allocas, so run
+ // even for optnone functions (skipFunction is intentionally not called).
bool Changed = false;
SmallVector<AllocaInst *, 8> Allocas;
for (auto &BB : F) {
diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
index 1186ae872d2c4..aba4d35e053f1 100644
--- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
@@ -11,6 +11,7 @@
//===----------------------------------------------------------------------===//
#include "NVPTXRegisterInfo.h"
+#include "MCTargetDesc/NVPTXBaseInfo.h"
#include "MCTargetDesc/NVPTXInstPrinter.h"
#include "NVPTX.h"
#include "NVPTXTargetMachine.h"
@@ -18,6 +19,7 @@
#include "llvm/CodeGen/MachineFrameInfo.h"
#include "llvm/CodeGen/MachineFunction.h"
#include "llvm/CodeGen/TargetInstrInfo.h"
+#include "llvm/IR/Instructions.h"
using namespace llvm;
@@ -113,11 +115,17 @@ bool NVPTXRegisterInfo::eliminateFrameIndex(MachineBasicBlock::iterator II,
const int FrameIndex = MI.getOperand(FIOperandNum).getIndex();
const MachineFunction &MF = *MI.getParent()->getParent();
- const int Offset = MF.getFrameInfo().getObjectOffset(FrameIndex) +
+ const MachineFrameInfo &MFI = MF.getFrameInfo();
+ const int Offset = MFI.getObjectOffset(FrameIndex) +
MI.getOperand(FIOperandNum + 1).getImm();
- // Using I0 as the frame pointer
- MI.getOperand(FIOperandNum).ChangeToRegister(getFrameRegister(MF), false);
+ // Local (addrspace 5) allocas are addressed through the local frame pointer
+ // (%SPL); everything else uses the generic frame pointer (%SP).
+ const AllocaInst *AI = MFI.getObjectAllocation(FrameIndex);
+ const Register FrameReg = AI && AI->getAddressSpace() == ADDRESS_SPACE_LOCAL
+ ? getFrameLocalRegister(MF)
+ : getFrameRegister(MF);
+ MI.getOperand(FIOperandNum).ChangeToRegister(FrameReg, false);
MI.getOperand(FIOperandNum + 1).ChangeToImmediate(Offset);
return false;
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
index fc954c3397caa..99b630c02b2c7 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
@@ -372,6 +372,10 @@ void NVPTXPassConfig::addIRPasses() {
if (getOptLevel() != CodeGenOptLevel::None) {
addAddressSpaceInferencePasses();
addStraightLineScalarOptimizationPasses();
+ } else {
+ // Required for correct stack lowering, not just an optimization, so run it
+ // even at -O0.
+ addPass(createNVPTXLowerAllocaPass());
}
addPass(createAtomicExpandLegacyPass());
diff --git a/llvm/test/CodeGen/NVPTX/call-with-alloca-buffer.ll b/llvm/test/CodeGen/NVPTX/call-with-alloca-buffer.ll
index 0eb7f6462f6fa..6a096365c2fd6 100644
--- a/llvm/test/CodeGen/NVPTX/call-with-alloca-buffer.ll
+++ b/llvm/test/CodeGen/NVPTX/call-with-alloca-buffer.ll
@@ -25,9 +25,10 @@ entry:
; CHECK: ld.param.b64 %rd[[A_REG:[0-9]+]], [kernel_func_param_0]
; CHECK: cvta.to.global.u64 %rd[[A1_REG:[0-9]+]], %rd[[A_REG]]
-; CHECK: add.u64 %rd[[SP_REG:[0-9]+]], %SP, 0
+; CHECK: add.u64 %rd[[BUF_REG:[0-9]+]], %SPL, 0
+; CHECK: cvta.local.u64 %rd[[SP_REG:[0-9]+]], %rd[[BUF_REG]]
; CHECK: ld.global.b32 %r[[A0_REG:[0-9]+]], [%rd[[A1_REG]]]
-; CHECK: st.local.b32 [{{%rd[0-9]+}}], %r[[A0_REG]]
+; CHECK: st.local.b32 [%SPL], %r[[A0_REG]]
%0 = load float, ptr %a, align 4
store float %0, ptr %buf, align 4
diff --git a/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll b/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
index 1510be1b0daa9..79be7ed2f18f1 100644
--- a/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
@@ -116,11 +116,11 @@ define %struct.64 @test_return_type_mismatch_variadic(ptr %p) {
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: mov.b64 %SPL, __local_depot3;
; CHECK-NEXT: ld.param.b64 %rd1, [test_return_type_mismatch_variadic_param_0];
-; CHECK-NEXT: add.u64 %rd2, %SPL, 0;
; CHECK-NEXT: { // callseq 3, 0
; CHECK-NEXT: .param .b64 param0;
; CHECK-NEXT: .param .b64 param1;
; CHECK-NEXT: .param .align 1 .b8 retval0[8];
+; CHECK-NEXT: add.u64 %rd2, %SPL, 0;
; CHECK-NEXT: st.param.b64 [param1], %rd2;
; CHECK-NEXT: st.param.b64 [param0], %rd1;
; CHECK-NEXT: mov.b64 %rd3, callee_variadic;
@@ -179,12 +179,12 @@ define i64 @test_param_type_mismatch_variadic(ptr %p) {
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: mov.b64 %SPL, __local_depot4;
; CHECK-NEXT: ld.param.b64 %rd1, [test_param_type_mismatch_variadic_param_0];
-; CHECK-NEXT: add.u64 %rd2, %SPL, 0;
-; CHECK-NEXT: st.local.b64 [%rd2], 7;
+; CHECK-NEXT: st.local.b64 [%SPL], 7;
; CHECK-NEXT: { // callseq 4, 0
; CHECK-NEXT: .param .b64 param0;
; CHECK-NEXT: .param .b64 param1;
; CHECK-NEXT: .param .b64 retval0;
+; CHECK-NEXT: add.u64 %rd2, %SPL, 0;
; CHECK-NEXT: st.param.b64 [param1], %rd2;
; CHECK-NEXT: st.param.b64 [param0], %rd1;
; CHECK-NEXT: call.uni (retval0), callee_variadic, (param0, param1);
@@ -207,12 +207,12 @@ define i64 @test_param_count_mismatch_variadic(ptr %p) {
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: mov.b64 %SPL, __local_depot5;
; CHECK-NEXT: ld.param.b64 %rd1, [test_param_count_mismatch_variadic_param_0];
-; CHECK-NEXT: add.u64 %rd2, %SPL, 0;
-; CHECK-NEXT: st.local.b64 [%rd2], 7;
+; CHECK-NEXT: st.local.b64 [%SPL], 7;
; CHECK-NEXT: { // callseq 5, 0
; CHECK-NEXT: .param .b64 param0;
; CHECK-NEXT: .param .b64 param1;
; CHECK-NEXT: .param .b64 retval0;
+; CHECK-NEXT: add.u64 %rd2, %SPL, 0;
; CHECK-NEXT: st.param.b64 [param1], %rd2;
; CHECK-NEXT: st.param.b64 [param0], %rd1;
; CHECK-NEXT: call.uni (retval0), callee_variadic, (param0, param1);
diff --git a/llvm/test/CodeGen/NVPTX/dynamic-stackalloc-regression.ll b/llvm/test/CodeGen/NVPTX/dynamic-stackalloc-regression.ll
index 0474d82556c1e..34702f1c177c5 100644
--- a/llvm/test/CodeGen/NVPTX/dynamic-stackalloc-regression.ll
+++ b/llvm/test/CodeGen/NVPTX/dynamic-stackalloc-regression.ll
@@ -13,13 +13,13 @@ define void @foo(i64 %a, ptr %p0, ptr %p1) {
; CHECK-NEXT: add.s64 %rd2, %rd1, 7;
; CHECK-NEXT: and.b64 %rd3, %rd2, -8;
; CHECK-NEXT: alloca.u64 %rd4, %rd3, 16;
-; CHECK-NEXT: cvta.local.u64 %rd5, %rd4;
-; CHECK-NEXT: ld.param.b64 %rd6, [foo_param_1];
-; CHECK-NEXT: alloca.u64 %rd7, %rd3, 16;
-; CHECK-NEXT: cvta.local.u64 %rd8, %rd7;
-; CHECK-NEXT: ld.param.b64 %rd9, [foo_param_2];
-; CHECK-NEXT: st.b64 [%rd6], %rd5;
-; CHECK-NEXT: st.b64 [%rd9], %rd8;
+; CHECK-NEXT: ld.param.b64 %rd5, [foo_param_1];
+; CHECK-NEXT: cvta.local.u64 %rd6, %rd4;
+; CHECK-NEXT: ld.param.b64 %rd7, [foo_param_2];
+; CHECK-NEXT: alloca.u64 %rd8, %rd3, 16;
+; CHECK-NEXT: cvta.local.u64 %rd9, %rd8;
+; CHECK-NEXT: st.b64 [%rd5], %rd6;
+; CHECK-NEXT: st.b64 [%rd7], %rd9;
; CHECK-NEXT: ret;
%b = alloca i8, i64 %a, align 16
%c = alloca i8, i64 %a, align 16
diff --git a/llvm/test/CodeGen/NVPTX/indirect_byval.ll b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
index c982aaf70d4a2..9b8f1f2f8346f 100644
--- a/llvm/test/CodeGen/NVPTX/indirect_byval.ll
+++ b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
@@ -21,18 +21,17 @@ define internal i32 @foo() {
; CHECK-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .align 1 .b8 _[1], .param .b64 _);
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: mov.b64 %SPL, __local_depot0;
-; CHECK-NEXT: cvta.local.u64 %SP, %SPL;
-; CHECK-NEXT: ld.global.b64 %rd1, [ptr];
+; CHECK-NEXT: add.u64 %rd1, %SPL, 0;
+; CHECK-NEXT: cvta.local.u64 %rd2, %rd1;
+; CHECK-NEXT: ld.global.b64 %rd3, [ptr];
; CHECK-NEXT: { // callseq 0, 0
; CHECK-NEXT: .param .align 1 .b8 param0[1];
; CHECK-NEXT: .param .b64 param1;
; CHECK-NEXT: .param .b32 retval0;
-; CHECK-NEXT: add.u64 %rd2, %SP, 0;
; CHECK-NEXT: st.param.b64 [param1], %rd2;
-; CHECK-NEXT: add.u64 %rd3, %SPL, 1;
-; CHECK-NEXT: ld.local.b8 %rs1, [%rd3];
+; CHECK-NEXT: ld.b8 %rs1, [%SPL+1];
; CHECK-NEXT: st.param.b8 [param0], %rs1;
-; CHECK-NEXT: call (retval0), %rd1, (param0, param1), prototype_0;
+; CHECK-NEXT: call (retval0), %rd3, (param0, param1), prototype_0;
; CHECK-NEXT: ld.param.b32 %r1, [retval0];
; CHECK-NEXT: } // callseq 0
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
@@ -57,18 +56,17 @@ define internal i32 @bar() {
; CHECK-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .align 8 .b8 _[8], .param .b64 _);
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: mov.b64 %SPL, __local_depot1;
-; CHECK-NEXT: cvta.local.u64 %SP, %SPL;
-; CHECK-NEXT: ld.global.b64 %rd1, [ptr];
+; CHECK-NEXT: add.u64 %rd1, %SPL, 0;
+; CHECK-NEXT: cvta.local.u64 %rd2, %rd1;
+; CHECK-NEXT: ld.global.b64 %rd3, [ptr];
; CHECK-NEXT: { // callseq 1, 0
; CHECK-NEXT: .param .align 8 .b8 param0[8];
; CHECK-NEXT: .param .b64 param1;
; CHECK-NEXT: .param .b32 retval0;
-; CHECK-NEXT: add.u64 %rd2, %SP, 0;
; CHECK-NEXT: st.param.b64 [param1], %rd2;
-; CHECK-NEXT: add.u64 %rd3, %SPL, 8;
-; CHECK-NEXT: ld.local.b64 %rd4, [%rd3];
+; CHECK-NEXT: ld.b64 %rd4, [%SPL+8];
; CHECK-NEXT: st.param.b64 [param0], %rd4;
-; CHECK-NEXT: call (retval0), %rd1, (param0, param1), prototype_1;
+; CHECK-NEXT: call (retval0), %rd3, (param0, param1), prototype_1;
; CHECK-NEXT: ld.param.b32 %r1, [retval0];
; CHECK-NEXT: } // callseq 1
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
diff --git a/llvm/test/CodeGen/NVPTX/local-stack-frame.ll b/llvm/test/CodeGen/NVPTX/local-stack-frame.ll
index 9dac46cb49005..0ce1653ecebcd 100644
--- a/llvm/test/CodeGen/NVPTX/local-stack-frame.ll
+++ b/llvm/test/CodeGen/NVPTX/local-stack-frame.ll
@@ -12,13 +12,12 @@ define void @foo(i32 %a) {
; PTX32-NEXT: .local .align 4 .b8 __local_depot0[4];
; PTX32-NEXT: .reg .b32 %SP;
; PTX32-NEXT: .reg .b32 %SPL;
-; PTX32-NEXT: .reg .b32 %r<3>;
+; PTX32-NEXT: .reg .b32 %r<2>;
; PTX32-EMPTY:
; PTX32-NEXT: // %bb.0:
; PTX32-NEXT: mov.b32 %SPL, __local_depot0;
; PTX32-NEXT: ld.param.b32 %r1, [foo_param_0];
-; PTX32-NEXT: add.u32 %r2, %SPL, 0;
-; PTX32-NEXT: st.local.b32 [%r2], %r1;
+; PTX32-NEXT: st.local.b32 [%SPL], %r1;
; PTX32-NEXT: ret;
;
; PTX64-LABEL: foo(
@@ -27,13 +26,11 @@ define void @foo(i32 %a) {
; PTX64-NEXT: .reg .b64 %SP;
; PTX64-NEXT: .reg .b64 %SPL;
; PTX64-NEXT: .reg .b32 %r<2>;
-; PTX64-NEXT: .reg .b64 %rd<2>;
; PTX64-EMPTY:
; PTX64-NEXT: // %bb.0:
; PTX64-NEXT: mov.b64 %SPL, __local_depot0;
; PTX64-NEXT: ld.param.b32 %r1, [foo_param_0];
-; PTX64-NEXT: add.u64 %rd1, %SPL, 0;
-; PTX64-NEXT: st.local.b32 [%rd1], %r1;
+; PTX64-NEXT: st.local.b32 [%SPL], %r1;
; PTX64-NEXT: ret;
%local = alloca i32, align 4
store volatile i32 %a, ptr %local
@@ -50,14 +47,13 @@ define ptx_kernel void @foo2(i32 %a) {
; PTX32-EMPTY:
; PTX32-NEXT: // %bb.0:
; PTX32-NEXT: mov.b32 %SPL, __local_depot1;
-; PTX32-NEXT: cvta.local.u32 %SP, %SPL;
; PTX32-NEXT: ld.param.b32 %r1, [foo2_param_0];
-; PTX32-NEXT: add.u32 %r2, %SP, 0;
-; PTX32-NEXT: add.u32 %r3, %SPL, 0;
-; PTX32-NEXT: st.local.b32 [%r3], %r1;
+; PTX32-NEXT: add.u32 %r2, %SPL, 0;
+; PTX32-NEXT: cvta.local.u32 %r3, %r2;
+; PTX32-NEXT: st.local.b32 [%SPL], %r1;
; PTX32-NEXT: { // callseq 0, 0
; PTX32-NEXT: .param .b32 param0;
-; PTX32-NEXT: st.param.b32 [param0], %r2;
+; PTX32-NEXT: st.param.b32 [param0], %r3;
; PTX32-NEXT: call.uni bar, (param0);
; PTX32-NEXT: } // callseq 0
; PTX32-NEXT: ret;
@@ -72,14 +68,13 @@ define ptx_kernel void @foo2(i32 %a) {
; PTX64-EMPTY:
; PTX64-NEXT: // %bb.0:
; PTX64-NEXT: mov.b64 %SPL, __local_depot1;
-; PTX64-NEXT: cvta.local.u64 %SP, %SPL;
; PTX64-NEXT: ld.param.b32 %r1, [foo2_param_0];
-; PTX64-NEXT: add.u64 %rd1, %SP, 0;
-; PTX64-NEXT: add.u64 %rd2, %SPL, 0;
-; PTX64-NEXT: st.local.b32 [%rd2], %r1;
+; PTX64-NEXT: add.u64 %rd1, %SPL, 0;
+; PTX64-NEXT: cvta.local.u64 %rd2, %rd1;
+; PTX64-NEXT: st.local.b32 [%SPL], %r1;
; PTX64-NEXT: { // callseq 0, 0
; PTX64-NEXT: .param .b64 param0;
-; PTX64-NEXT: st.param.b64 [param0], %rd1;
+; PTX64-NEXT: st.param.b64 [param0], %rd2;
; PTX64-NEXT: call.uni bar, (param0);
; PTX64-NEXT: } // callseq 0
; PTX64-NEXT: ret;
@@ -102,9 +97,9 @@ define void @foo3(i32 %a) {
; PTX32-NEXT: // %bb.0:
; PTX32-NEXT: mov.b32 %SPL, __local_depot2;
; PTX32-NEXT: ld.param.b32 %r1, [foo3_param_0];
-; PTX32-NEXT: add.u32 %r2, %SPL, 0;
-; PTX32-NEXT: shl.b32 %r3, %r1, 2;
-; PTX32-NEXT: add.s32 %r4, %r2, %r3;
+; PTX32-NEXT: shl.b32 %r2, %r1, 2;
+; PTX32-NEXT: add.u32 %r3, %SPL, 0;
+; PTX32-NEXT: add.s32 %r4, %r3, %r2;
; PTX32-NEXT: st.local.b32 [%r4], %r1;
; PTX32-NEXT: ret;
;
@@ -119,9 +114,9 @@ define void @foo3(i32 %a) {
; PTX64-NEXT: // %bb.0:
; PTX64-NEXT: mov.b64 %SPL, __local_depot2;
; PTX64-NEXT: ld.param.b32 %r1, [foo3_param_0];
-; PTX64-NEXT: add.u64 %rd1, %SPL, 0;
-; PTX64-NEXT: mul.wide.s32 %rd2, %r1, 4;
-; PTX64-NEXT: add.s64 %rd3, %rd1, %rd2;
+; PTX64-NEXT: mul.wide.s32 %rd1, %r1, 4;
+; PTX64-NEXT: add.u64 %rd2, %SPL, 0;
+; PTX64-NEXT: add.s64 %rd3, %rd2, %rd1;
; PTX64-NEXT: st.local.b32 [%rd3], %r1;
; PTX64-NEXT: ret;
%local = alloca [3 x i32], align 4
@@ -140,21 +135,20 @@ define void @foo4() {
; PTX32-EMPTY:
; PTX32-NEXT: // %bb.0:
; PTX32-NEXT: mov.b32 %SPL, __local_depot3;
-; PTX32-NEXT: cvta.local.u32 %SP, %SPL;
-; PTX32-NEXT: add.u32 %r1, %SP, 0;
-; PTX32-NEXT: add.u32 %r2, %SPL, 0;
-; PTX32-NEXT: add.u32 %r3, %SP, 4;
-; PTX32-NEXT: add.u32 %r4, %SPL, 4;
-; PTX32-NEXT: st.local.b32 [%r2], 0;
-; PTX32-NEXT: st.local.b32 [%r4], 0;
+; PTX32-NEXT: add.u32 %r1, %SPL, 0;
+; PTX32-NEXT: cvta.local.u32 %r2, %r1;
+; PTX32-NEXT: add.u32 %r3, %SPL, 4;
+; PTX32-NEXT: cvta.local.u32 %r4, %r3;
+; PTX32-NEXT: st.local.b32 [%SPL], 0;
+; PTX32-NEXT: st.local.b32 [%SPL+4], 0;
; PTX32-NEXT: { // callseq 1, 0
; PTX32-NEXT: .param .b32 param0;
-; PTX32-NEXT: st.param.b32 [param0], %r1;
+; PTX32-NEXT: st.param.b32 [param0], %r2;
; PTX32-NEXT: call.uni bar, (param0);
; PTX32-NEXT: } // callseq 1
; PTX32-NEXT: { // callseq 2, 0
; PTX32-NEXT: .param .b32 param0;
-; PTX32-NEXT: st.param.b32 [param0], %r3;
+; PTX32-NEXT: st.param.b32 [param0], %r4;
; PTX32-NEXT: call.uni bar, (param0);
; PTX32-NEXT: } // callseq 2
; PTX32-NEXT: ret;
@@ -168,21 +162,20 @@ define void @foo4() {
; PTX64-EMPTY:
; PTX64-NEXT: // %bb.0:
; PTX64-NEXT: mov.b64 %SPL, __local_depot3;
-; PTX64-NEXT: cvta.local.u64 %SP, %SPL;
-; PTX64-NEXT: add.u64 %rd1, %SP, 0;
-; PTX64-NEXT: add.u64 %rd2, %SPL, 0;
-; PTX64-NEXT: add.u64 %rd3, %SP, 4;
-; PTX64-NEXT: add.u64 %rd4, %SPL, 4;
-; PTX64-NEXT: st.local.b32 [%rd2], 0;
-; PTX64-NEXT: st.local.b32 [%rd4], 0;
+; PTX64-NEXT: add.u64 %rd1, %SPL, 0;
+; PTX64-NEXT: cvta.local.u64 %rd2, %rd1;
+; PTX64-NEXT: add.u64 %rd3, %SPL, 4;
+; PTX64-NEXT: cvta.local.u64 %rd4, %rd3;
+; PTX64-NEXT: st.local.b32 [%SPL], 0;
+; PTX64-NEXT: st.local.b32 [%SPL+4], 0;
; PTX64-NEXT: { // callseq 1, 0
; PTX64-NEXT: .param .b64 param0;
-; PTX64-NEXT: st.param.b64 [param0], %rd1;
+; PTX64-NEXT: st.param.b64 [param0], %rd2;
; PTX64-NEXT: call.uni bar, (param0);
; PTX64-NEXT: } // callseq 1
; PTX64-NEXT: { // callseq 2, 0
; PTX64-NEXT: .param .b64 param0;
-; PTX64-NEXT: st.param.b64 [param0], %rd3;
+; PTX64-NEXT: st.param.b64 [param0], %rd4;
; PTX64-NEXT: call.uni bar, (param0);
; PTX64-NEXT: } // callseq 2
; PTX64-NEXT: ret;
diff --git a/llvm/test/CodeGen/NVPTX/lower-alloca.ll b/llvm/test/CodeGen/NVPTX/lower-alloca.ll
index d95619d5c842a..b93e6a316ddf4 100644
--- a/llvm/test/CodeGen/NVPTX/lower-alloca.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-alloca.ll
@@ -18,7 +18,7 @@ define ptx_kernel void @kernel() {
; LOWERALLOCAONLY: [[GENERIC:%.*]] = addrspacecast ptr addrspace(5) [[A]] to ptr
; LOWERALLOCAONLY: store i32 0, ptr addrspace(5) [[A]], align 4
; LOWERALLOCAONLY: call void @callee(ptr [[GENERIC]])
-; PTX: st.local.b32 [{{.*}}], 0
+; PTX: st.local.b32 [%SPL], 0
store i32 0, ptr %A
call void @callee(ptr %A)
ret void
@@ -29,7 +29,7 @@ define void @alloca_in_explicit_local_as() {
; PTX-LABEL: .visible .func alloca_in_explicit_local_as(
%A = alloca i32, addrspace(5)
; CHECK: store i32 0, ptr addrspace(5) {{%.+}}
-; PTX: st.local.b32 [%SP], 0
+; PTX: st.local.b32 [%SPL], 0
; LOWERALLOCAONLY: [[V1:%.*]] = addrspacecast ptr addrspace(5) %A to ptr
; LOWERALLOCAONLY: store i32 0, ptr [[V1]], align 4
store i32 0, ptr addrspace(5) %A
diff --git a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
index 6700fa3c3f4e6..57c9abe372f48 100644
--- a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
@@ -147,14 +147,13 @@ define dso_local ptx_kernel void @escape_ptr(ptr nocapture noundef readnone %out
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %entry
; PTX-NEXT: mov.b64 %SPL, __local_depot2;
-; PTX-NEXT: cvta.local.u64 %SP, %SPL;
-; PTX-NEXT: add.u64 %rd1, %SP, 0;
+; PTX-NEXT: add.u64 %rd1, %SPL, 0;
; PTX-NEXT: cvta.local.u64 %rd2, %rd1;
; PTX-NEXT: ld.param.b32 %rd3, [escape_ptr_param_1+4];
; PTX-NEXT: shl.b64 %rd4, %rd3, 32;
; PTX-NEXT: ld.param.b32 %rd5, [escape_ptr_param_1];
; PTX-NEXT: or.b64 %rd6, %rd4, %rd5;
-; PTX-NEXT: st.local.b64 [%SP], %rd6;
+; PTX-NEXT: st.local.b64 [%SPL], %rd6;
; PTX-NEXT: { // callseq 0, 0
; PTX-NEXT: .param .b64 param0;
; PTX-NEXT: st.param.b64 [param0], %rd2;
@@ -195,13 +194,12 @@ define dso_local ptx_kernel void @escape_ptr_gep(ptr nocapture noundef readnone
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %entry
; PTX-NEXT: mov.b64 %SPL, __local_depot3;
-; PTX-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-NEXT: ld.param.b32 %rd1, [escape_ptr_gep_param_1+4];
; PTX-NEXT: shl.b64 %rd2, %rd1, 32;
; PTX-NEXT: ld.param.b32 %rd3, [escape_ptr_gep_param_1];
; PTX-NEXT: or.b64 %rd4, %rd2, %rd3;
-; PTX-NEXT: st.local.b64 [%SP], %rd4;
-; PTX-NEXT: add.u64 %rd5, %SP, 0;
+; PTX-NEXT: st.local.b64 [%SPL], %rd4;
+; PTX-NEXT: add.u64 %rd5, %SPL, 0;
; PTX-NEXT: or.b64 %rd6, %rd5, 4;
; PTX-NEXT: cvta.local.u64 %rd7, %rd6;
; PTX-NEXT: { // callseq 1, 0
@@ -243,16 +241,15 @@ define dso_local ptx_kernel void @escape_ptr_store(ptr nocapture noundef writeon
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %entry
; PTX-NEXT: mov.b64 %SPL, __local_depot4;
-; PTX-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-NEXT: ld.param.b64 %rd1, [escape_ptr_store_param_0];
; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;
-; PTX-NEXT: add.u64 %rd3, %SP, 0;
+; PTX-NEXT: add.u64 %rd3, %SPL, 0;
; PTX-NEXT: cvta.local.u64 %rd4, %rd3;
; PTX-NEXT: ld.param.b32 %rd5, [escape_ptr_store_param_1+4];
; PTX-NEXT: shl.b64 %rd6, %rd5, 32;
; PTX-NEXT: ld.param.b32 %rd7, [escape_ptr_store_param_1];
; PTX-NEXT: or.b64 %rd8, %rd6, %rd7;
-; PTX-NEXT: st.local.b64 [%SP], %rd8;
+; PTX-NEXT: st.local.b64 [%SPL], %rd8;
; PTX-NEXT: st.global.b64 [%rd2], %rd4;
; PTX-NEXT: ret;
entry:
@@ -289,15 +286,14 @@ define dso_local ptx_kernel void @escape_ptr_gep_store(ptr nocapture noundef wri
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %entry
; PTX-NEXT: mov.b64 %SPL, __local_depot5;
-; PTX-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-NEXT: ld.param.b64 %rd1, [escape_ptr_gep_store_param_0];
; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;
; PTX-NEXT: ld.param.b32 %rd3, [escape_ptr_gep_store_param_1+4];
; PTX-NEXT: shl.b64 %rd4, %rd3, 32;
; PTX-NEXT: ld.param.b32 %rd5, [escape_ptr_gep_store_param_1];
; PTX-NEXT: or.b64 %rd6, %rd4, %rd5;
-; PTX-NEXT: st.local.b64 [%SP], %rd6;
-; PTX-NEXT: add.u64 %rd7, %SP, 0;
+; PTX-NEXT: st.local.b64 [%SPL], %rd6;
+; PTX-NEXT: add.u64 %rd7, %SPL, 0;
; PTX-NEXT: or.b64 %rd8, %rd7, 4;
; PTX-NEXT: cvta.local.u64 %rd9, %rd8;
; PTX-NEXT: st.global.b64 [%rd2], %rd9;
@@ -337,16 +333,15 @@ define dso_local ptx_kernel void @escape_ptrtoint(ptr nocapture noundef writeonl
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %entry
; PTX-NEXT: mov.b64 %SPL, __local_depot6;
-; PTX-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-NEXT: ld.param.b64 %rd1, [escape_ptrtoint_param_0];
; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;
-; PTX-NEXT: add.u64 %rd3, %SP, 0;
+; PTX-NEXT: add.u64 %rd3, %SPL, 0;
; PTX-NEXT: cvta.local.u64 %rd4, %rd3;
; PTX-NEXT: ld.param.b32 %rd5, [escape_ptrtoint_param_1+4];
; PTX-NEXT: shl.b64 %rd6, %rd5, 32;
; PTX-NEXT: ld.param.b32 %rd7, [escape_ptrtoint_param_1];
; PTX-NEXT: or.b64 %rd8, %rd6, %rd7;
-; PTX-NEXT: st.local.b64 [%SP], %rd8;
+; PTX-NEXT: st.local.b64 [%SPL], %rd8;
; PTX-NEXT: st.global.b64 [%rd2], %rd4;
; PTX-NEXT: ret;
entry:
@@ -501,12 +496,11 @@ define dso_local ptx_kernel void @memcpy_to_param(ptr nocapture noundef readonly
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %entry
; PTX-NEXT: mov.b64 %SPL, __local_depot9;
-; PTX-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-NEXT: ld.param.b64 %rd1, [memcpy_to_param_param_0];
; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;
; PTX-NEXT: ld.param.b32 %r1, [memcpy_to_param_param_1+4];
; PTX-NEXT: ld.param.b32 %r2, [memcpy_to_param_param_1];
-; PTX-NEXT: st.local.v2.b32 [%SP], {%r2, %r1};
+; PTX-NEXT: st.local.v2.b32 [%SPL], {%r2, %r1};
; PTX-NEXT: ld.volatile.global.b8 %r3, [%rd2+4];
; PTX-NEXT: ld.volatile.global.b8 %r4, [%rd2+5];
; PTX-NEXT: shl.b32 %r5, %r4, 8;
@@ -527,7 +521,7 @@ define dso_local ptx_kernel void @memcpy_to_param(ptr nocapture noundef readonly
; PTX-NEXT: shl.b32 %r20, %r19, 24;
; PTX-NEXT: or.b32 %r21, %r20, %r18;
; PTX-NEXT: or.b32 %r22, %r21, %r16;
-; PTX-NEXT: st.local.v2.b32 [%SP], {%r22, %r12};
+; PTX-NEXT: st.local.v2.b32 [%SPL], {%r22, %r12};
; PTX-NEXT: ret;
entry:
tail call void @llvm.memcpy.p0.p0.i64(ptr %s, ptr %in, i64 16, i1 true)
@@ -662,16 +656,15 @@ define ptx_kernel void @test_select_write(ptr byval(i32) align 4 %input1, ptr by
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %bb
; PTX-NEXT: mov.b64 %SPL, __local_depot12;
-; PTX-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-NEXT: ld.param.b8 %rs1, [test_select_write_param_3];
; PTX-NEXT: and.b16 %rs2, %rs1, 1;
; PTX-NEXT: setp.ne.b16 %p1, %rs2, 0;
; PTX-NEXT: ld.param.b32 %r1, [test_select_write_param_1];
-; PTX-NEXT: st.local.b32 [%SP], %r1;
+; PTX-NEXT: st.local.b32 [%SPL], %r1;
; PTX-NEXT: ld.param.b32 %r2, [test_select_write_param_0];
-; PTX-NEXT: st.local.b32 [%SP+4], %r2;
-; PTX-NEXT: add.u64 %rd1, %SP, 0;
-; PTX-NEXT: add.u64 %rd2, %SP, 4;
+; PTX-NEXT: st.local.b32 [%SPL+4], %r2;
+; PTX-NEXT: add.u64 %rd1, %SPL, 0;
+; PTX-NEXT: add.u64 %rd2, %SPL, 4;
; PTX-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1;
; PTX-NEXT: st.local.b32 [%rd3], 1;
; PTX-NEXT: ret;
@@ -852,18 +845,17 @@ define ptx_kernel void @test_phi_write(ptr byval(%struct.S) align 4 %input1, ptr
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %bb
; PTX-NEXT: mov.b64 %SPL, __local_depot14;
-; PTX-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-NEXT: ld.param.b8 %rs1, [test_phi_write_param_2];
; PTX-NEXT: and.b16 %rs2, %rs1, 1;
; PTX-NEXT: setp.ne.b16 %p1, %rs2, 0;
; PTX-NEXT: ld.param.b32 %r1, [test_phi_write_param_1+4];
-; PTX-NEXT: st.local.b32 [%SP], %r1;
+; PTX-NEXT: st.local.b32 [%SPL], %r1;
; PTX-NEXT: ld.param.b32 %r2, [test_phi_write_param_0];
-; PTX-NEXT: st.local.b32 [%SP+4], %r2;
-; PTX-NEXT: add.u64 %rd1, %SP, 4;
+; PTX-NEXT: st.local.b32 [%SPL+4], %r2;
+; PTX-NEXT: add.u64 %rd1, %SPL, 4;
; PTX-NEXT: @%p1 bra $L__BB14_2;
; PTX-NEXT: // %bb.1: // %second
-; PTX-NEXT: add.u64 %rd1, %SP, 0;
+; PTX-NEXT: add.u64 %rd1, %SPL, 0;
; PTX-NEXT: $L__BB14_2: // %merge
; PTX-NEXT: st.local.b32 [%rd1], 1;
; PTX-NEXT: ret;
@@ -908,12 +900,11 @@ define ptx_kernel void @test_forward_byval_arg(ptr byval(i32) align 4 %input) {
; PTX-EMPTY:
; PTX-NEXT: // %bb.0:
; PTX-NEXT: mov.b64 %SPL, __local_depot15;
-; PTX-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-NEXT: ld.param.b32 %r1, [test_forward_byval_arg_param_0];
-; PTX-NEXT: st.local.b32 [%SP], %r1;
+; PTX-NEXT: st.local.b32 [%SPL], %r1;
; PTX-NEXT: { // callseq 2, 0
; PTX-NEXT: .param .align 4 .b8 param0[4];
-; PTX-NEXT: ld.b32 %r2, [%SP];
+; PTX-NEXT: ld.b32 %r2, [%SPL];
; PTX-NEXT: st.param.b32 [param0], %r2;
; PTX-NEXT: call.uni device_func, (param0);
; PTX-NEXT: } // callseq 2
diff --git a/llvm/test/CodeGen/NVPTX/vaargs.ll b/llvm/test/CodeGen/NVPTX/vaargs.ll
index ea0807a88bd68..8802ae3be1816 100644
--- a/llvm/test/CodeGen/NVPTX/vaargs.ll
+++ b/llvm/test/CodeGen/NVPTX/vaargs.ll
@@ -21,36 +21,35 @@ define i32 @foo(i32 %a, ...) {
; CHECK32-EMPTY:
; CHECK32-NEXT: // %bb.0: // %entry
; CHECK32-NEXT: mov.b32 %SPL, __local_depot0;
-; CHECK32-NEXT: cvta.local.u32 %SP, %SPL;
; CHECK32-NEXT: ld.param.b32 %r2, [foo_param_1];
; CHECK32-NEXT: ld.param.b32 %r1, [foo_param_0];
; CHECK32-NEXT: cvta.local.u32 %r3, %r2;
-; CHECK32-NEXT: st.b32 [%SP], %r3;
-; CHECK32-NEXT: ld.b32 %r4, [%SP];
-; CHECK32-NEXT: st.b32 [%SP+16], %r4;
-; CHECK32-NEXT: ld.b32 %r5, [%SP];
+; CHECK32-NEXT: st.local.b32 [%SPL], %r3;
+; CHECK32-NEXT: ld.local.b32 %r4, [%SPL];
+; CHECK32-NEXT: st.local.b32 [%SPL+16], %r4;
+; CHECK32-NEXT: ld.local.b32 %r5, [%SPL];
; CHECK32-NEXT: add.s32 %r6, %r5, 3;
; CHECK32-NEXT: and.b32 %r7, %r6, -4;
; CHECK32-NEXT: add.s32 %r8, %r7, 4;
-; CHECK32-NEXT: st.b32 [%SP], %r8;
+; CHECK32-NEXT: st.local.b32 [%SPL], %r8;
; CHECK32-NEXT: ld.b32 %r9, [%r7];
-; CHECK32-NEXT: ld.b32 %r10, [%SP];
+; CHECK32-NEXT: ld.local.b32 %r10, [%SPL];
; CHECK32-NEXT: add.s32 %r11, %r10, 7;
; CHECK32-NEXT: and.b32 %r12, %r11, -8;
; CHECK32-NEXT: add.s32 %r13, %r12, 8;
-; CHECK32-NEXT: st.b32 [%SP], %r13;
+; CHECK32-NEXT: st.local.b32 [%SPL], %r13;
; CHECK32-NEXT: ld.b64 %rd1, [%r12];
-; CHECK32-NEXT: ld.b32 %r14, [%SP];
+; CHECK32-NEXT: ld.local.b32 %r14, [%SPL];
; CHECK32-NEXT: add.s32 %r15, %r14, 7;
; CHECK32-NEXT: and.b32 %r16, %r15, -8;
; CHECK32-NEXT: add.s32 %r17, %r16, 8;
-; CHECK32-NEXT: st.b32 [%SP], %r17;
+; CHECK32-NEXT: st.local.b32 [%SPL], %r17;
; CHECK32-NEXT: ld.b64 %rd2, [%r16];
-; CHECK32-NEXT: ld.b32 %r18, [%SP];
+; CHECK32-NEXT: ld.local.b32 %r18, [%SPL];
; CHECK32-NEXT: add.s32 %r19, %r18, 3;
; CHECK32-NEXT: and.b32 %r20, %r19, -4;
; CHECK32-NEXT: add.s32 %r21, %r20, 4;
-; CHECK32-NEXT: st.b32 [%SP], %r21;
+; CHECK32-NEXT: st.local.b32 [%SPL], %r21;
; CHECK32-NEXT: ld.b32 %r22, [%r20];
; CHECK32-NEXT: { // callseq 0, 0
; CHECK32-NEXT: .param .b32 param0;
@@ -67,11 +66,11 @@ define i32 @foo(i32 %a, ...) {
; CHECK32-NEXT: call.uni (retval0), bar, (param0, param1, param2, param3, param4);
; CHECK32-NEXT: ld.param.b32 %r23, [retval0];
; CHECK32-NEXT: } // callseq 0
-; CHECK32-NEXT: ld.b32 %r24, [%SP+16];
+; CHECK32-NEXT: ld.local.b32 %r24, [%SPL+16];
; CHECK32-NEXT: add.s32 %r25, %r24, 3;
; CHECK32-NEXT: and.b32 %r26, %r25, -4;
; CHECK32-NEXT: add.s32 %r27, %r26, 4;
-; CHECK32-NEXT: st.b32 [%SP+16], %r27;
+; CHECK32-NEXT: st.local.b32 [%SPL+16], %r27;
; CHECK32-NEXT: ld.b32 %r28, [%r26];
; CHECK32-NEXT: add.s32 %r29, %r23, %r28;
; CHECK32-NEXT: st.param.b32 [func_retval0], %r29;
@@ -87,36 +86,35 @@ define i32 @foo(i32 %a, ...) {
; CHECK64-EMPTY:
; CHECK64-NEXT: // %bb.0: // %entry
; CHECK64-NEXT: mov.b64 %SPL, __local_depot0;
-; CHECK64-NEXT: cvta.local.u64 %SP, %SPL;
; CHECK64-NEXT: ld.param.b64 %rd1, [foo_param_1];
; CHECK64-NEXT: ld.param.b32 %r1, [foo_param_0];
; CHECK64-NEXT: cvta.local.u64 %rd2, %rd1;
-; CHECK64-NEXT: st.b64 [%SP], %rd2;
-; CHECK64-NEXT: ld.b64 %rd3, [%SP];
-; CHECK64-NEXT: st.b64 [%SP+24], %rd3;
-; CHECK64-NEXT: ld.b64 %rd4, [%SP];
+; CHECK64-NEXT: st.local.b64 [%SPL], %rd2;
+; CHECK64-NEXT: ld.local.b64 %rd3, [%SPL];
+; CHECK64-NEXT: st.local.b64 [%SPL+24], %rd3;
+; CHECK64-NEXT: ld.local.b64 %rd4, [%SPL];
; CHECK64-NEXT: add.s64 %rd5, %rd4, 3;
; CHECK64-NEXT: and.b64 %rd6, %rd5, -4;
; CHECK64-NEXT: add.s64 %rd7, %rd6, 4;
-; CHECK64-NEXT: st.b64 [%SP], %rd7;
+; CHECK64-NEXT: st.local.b64 [%SPL], %rd7;
; CHECK64-NEXT: ld.b32 %r2, [%rd6];
-; CHECK64-NEXT: ld.b64 %rd8, [%SP];
+; CHECK64-NEXT: ld.local.b64 %rd8, [%SPL];
; CHECK64-NEXT: add.s64 %rd9, %rd8, 7;
; CHECK64-NEXT: and.b64 %rd10, %rd9, -8;
; CHECK64-NEXT: add.s64 %rd11, %rd10, 8;
-; CHECK64-NEXT: st.b64 [%SP], %rd11;
+; CHECK64-NEXT: st.local.b64 [%SPL], %rd11;
; CHECK64-NEXT: ld.b64 %rd12, [%rd10];
-; CHECK64-NEXT: ld.b64 %rd13, [%SP];
+; CHECK64-NEXT: ld.local.b64 %rd13, [%SPL];
; CHECK64-NEXT: add.s64 %rd14, %rd13, 7;
; CHECK64-NEXT: and.b64 %rd15, %rd14, -8;
; CHECK64-NEXT: add.s64 %rd16, %rd15, 8;
-; CHECK64-NEXT: st.b64 [%SP], %rd16;
+; CHECK64-NEXT: st.local.b64 [%SPL], %rd16;
; CHECK64-NEXT: ld.b64 %rd17, [%rd15];
-; CHECK64-NEXT: ld.b64 %rd18, [%SP];
+; CHECK64-NEXT: ld.local.b64 %rd18, [%SPL];
; CHECK64-NEXT: add.s64 %rd19, %rd18, 7;
; CHECK64-NEXT: and.b64 %rd20, %rd19, -8;
; CHECK64-NEXT: add.s64 %rd21, %rd20, 8;
-; CHECK64-NEXT: st.b64 [%SP], %rd21;
+; CHECK64-NEXT: st.local.b64 [%SPL], %rd21;
; CHECK64-NEXT: ld.b64 %rd22, [%rd20];
; CHECK64-NEXT: { // callseq 0, 0
; CHECK64-NEXT: .param .b32 param0;
@@ -133,11 +131,11 @@ define i32 @foo(i32 %a, ...) {
; CHECK64-NEXT: call.uni (retval0), bar, (param0, param1, param2, param3, param4);
; CHECK64-NEXT: ld.param.b32 %r3, [retval0];
; CHECK64-NEXT: } // callseq 0
-; CHECK64-NEXT: ld.b64 %rd23, [%SP+24];
+; CHECK64-NEXT: ld.local.b64 %rd23, [%SPL+24];
; CHECK64-NEXT: add.s64 %rd24, %rd23, 3;
; CHECK64-NEXT: and.b64 %rd25, %rd24, -4;
; CHECK64-NEXT: add.s64 %rd26, %rd25, 4;
-; CHECK64-NEXT: st.b64 [%SP+24], %rd26;
+; CHECK64-NEXT: st.local.b64 [%SPL+24], %rd26;
; CHECK64-NEXT: ld.b32 %r4, [%rd25];
; CHECK64-NEXT: add.s32 %r5, %r3, %r4;
; CHECK64-NEXT: st.param.b32 [func_retval0], %r5;
@@ -186,12 +184,11 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK32-NEXT: .local .align 8 .b8 __local_depot1[32];
; CHECK32-NEXT: .reg .b32 %SP;
; CHECK32-NEXT: .reg .b32 %SPL;
-; CHECK32-NEXT: .reg .b32 %r<9>;
+; CHECK32-NEXT: .reg .b32 %r<8>;
; CHECK32-NEXT: .reg .b64 %rd<3>;
-; CHECK32-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _);
+; CHECK32-EMPTY:
; CHECK32-NEXT: // %bb.0: // %entry
; CHECK32-NEXT: mov.b32 %SPL, __local_depot1;
-; CHECK32-NEXT: cvta.local.u32 %SP, %SPL;
; CHECK32-NEXT: ld.param.b32 %r2, [test_foo_param_3];
; CHECK32-NEXT: ld.param.b64 %rd2, [test_foo_param_2];
; CHECK32-NEXT: ld.param.b64 %rd1, [test_foo_param_1];
@@ -199,22 +196,22 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK32-NEXT: mov.b32 %r3, foo_ptr;
; CHECK32-NEXT: cvta.global.u32 %r4, %r3;
; CHECK32-NEXT: ld.b32 %r5, [%r4];
-; CHECK32-NEXT: st.b32 [%SP], %r1;
-; CHECK32-NEXT: st.b64 [%SP+8], %rd1;
-; CHECK32-NEXT: st.b64 [%SP+16], %rd2;
-; CHECK32-NEXT: st.b32 [%SP+24], %r2;
-; CHECK32-NEXT: add.u32 %r6, %SP, 0;
-; CHECK32-NEXT: cvta.to.local.u32 %r7, %r6;
+; CHECK32-NEXT: st.local.b32 [%SPL], %r1;
+; CHECK32-NEXT: st.local.b64 [%SPL+8], %rd1;
+; CHECK32-NEXT: st.local.b64 [%SPL+16], %rd2;
+; CHECK32-NEXT: st.local.b32 [%SPL+24], %r2;
; CHECK32-NEXT: { // callseq 1, 0
; CHECK32-NEXT: .param .b32 param0;
; CHECK32-NEXT: .param .b32 param1;
; CHECK32-NEXT: .param .b32 retval0;
-; CHECK32-NEXT: st.param.b32 [param1], %r7;
+; CHECK32-NEXT: add.u32 %r6, %SPL, 0;
+; CHECK32-NEXT: st.param.b32 [param1], %r6;
+; CHECK32-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _);
; CHECK32-NEXT: st.param.b32 [param0], 4;
; CHECK32-NEXT: call (retval0), %r5, (param0, param1), prototype_1;
-; CHECK32-NEXT: ld.param.b32 %r8, [retval0];
+; CHECK32-NEXT: ld.param.b32 %r7, [retval0];
; CHECK32-NEXT: } // callseq 1
-; CHECK32-NEXT: st.param.b32 [func_retval0], %r8;
+; CHECK32-NEXT: st.param.b32 [func_retval0], %r7;
; CHECK32-NEXT: ret;
;
; CHECK64-LABEL: test_foo(
@@ -223,11 +220,10 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK64-NEXT: .reg .b64 %SP;
; CHECK64-NEXT: .reg .b64 %SPL;
; CHECK64-NEXT: .reg .b32 %r<3>;
-; CHECK64-NEXT: .reg .b64 %rd<9>;
-; CHECK64-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b64 _);
+; CHECK64-NEXT: .reg .b64 %rd<8>;
+; CHECK64-EMPTY:
; CHECK64-NEXT: // %bb.0: // %entry
; CHECK64-NEXT: mov.b64 %SPL, __local_depot1;
-; CHECK64-NEXT: cvta.local.u64 %SP, %SPL;
; CHECK64-NEXT: ld.param.b64 %rd3, [test_foo_param_3];
; CHECK64-NEXT: ld.param.b64 %rd2, [test_foo_param_2];
; CHECK64-NEXT: ld.param.b64 %rd1, [test_foo_param_1];
@@ -235,17 +231,17 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK64-NEXT: mov.b64 %rd4, foo_ptr;
; CHECK64-NEXT: cvta.global.u64 %rd5, %rd4;
; CHECK64-NEXT: ld.b64 %rd6, [%rd5];
-; CHECK64-NEXT: st.b32 [%SP], %r1;
-; CHECK64-NEXT: st.b64 [%SP+8], %rd1;
-; CHECK64-NEXT: st.b64 [%SP+16], %rd2;
-; CHECK64-NEXT: st.b64 [%SP+24], %rd3;
-; CHECK64-NEXT: add.u64 %rd7, %SP, 0;
-; CHECK64-NEXT: cvta.to.local.u64 %rd8, %rd7;
+; CHECK64-NEXT: st.local.b32 [%SPL], %r1;
+; CHECK64-NEXT: st.local.b64 [%SPL+8], %rd1;
+; CHECK64-NEXT: st.local.b64 [%SPL+16], %rd2;
+; CHECK64-NEXT: st.local.b64 [%SPL+24], %rd3;
; CHECK64-NEXT: { // callseq 1, 0
; CHECK64-NEXT: .param .b32 param0;
; CHECK64-NEXT: .param .b64 param1;
; CHECK64-NEXT: .param .b32 retval0;
-; CHECK64-NEXT: st.param.b64 [param1], %rd8;
+; CHECK64-NEXT: add.u64 %rd7, %SPL, 0;
+; CHECK64-NEXT: st.param.b64 [param1], %rd7;
+; CHECK64-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b64 _);
; CHECK64-NEXT: st.param.b32 [param0], 4;
; CHECK64-NEXT: call (retval0), %rd6, (param0, param1), prototype_1;
; CHECK64-NEXT: ld.param.b32 %r2, [retval0];
diff --git a/llvm/test/CodeGen/NVPTX/variadics-backend.ll b/llvm/test/CodeGen/NVPTX/variadics-backend.ll
index 181d38093b0ba..e010e71459bca 100644
--- a/llvm/test/CodeGen/NVPTX/variadics-backend.ll
+++ b/llvm/test/CodeGen/NVPTX/variadics-backend.ll
@@ -109,17 +109,16 @@ define dso_local i32 @foo() {
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot1;
-; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;
-; CHECK-PTX-NEXT: st.local.b64 [%SP], 4294967297;
-; CHECK-PTX-NEXT: st.local.b32 [%SP+8], 1;
-; CHECK-PTX-NEXT: st.local.b64 [%SP+16], 1;
-; CHECK-PTX-NEXT: st.local.b64 [%SP+24], 4607182418800017408;
-; CHECK-PTX-NEXT: st.local.b64 [%SP+32], 4607182418800017408;
+; CHECK-PTX-NEXT: st.local.b64 [%SPL], 4294967297;
+; CHECK-PTX-NEXT: st.local.b32 [%SPL+8], 1;
+; CHECK-PTX-NEXT: st.local.b64 [%SPL+16], 1;
+; CHECK-PTX-NEXT: st.local.b64 [%SPL+24], 4607182418800017408;
+; CHECK-PTX-NEXT: st.local.b64 [%SPL+32], 4607182418800017408;
; CHECK-PTX-NEXT: { // callseq 0, 0
; CHECK-PTX-NEXT: .param .b32 param0;
; CHECK-PTX-NEXT: .param .b64 param1;
; CHECK-PTX-NEXT: .param .b32 retval0;
-; CHECK-PTX-NEXT: add.u64 %rd1, %SP, 0;
+; CHECK-PTX-NEXT: add.u64 %rd1, %SPL, 0;
; CHECK-PTX-NEXT: st.param.b64 [param1], %rd1;
; CHECK-PTX-NEXT: st.param.b32 [param0], 1;
; CHECK-PTX-NEXT: call.uni (retval0), variadics1, (param0, param1);
@@ -147,7 +146,6 @@ define dso_local i32 @variadics2(i32 noundef %first, ...) {
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot2;
-; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;
; CHECK-PTX-NEXT: ld.param.b32 %r1, [variadics2_param_0];
; CHECK-PTX-NEXT: ld.param.b64 %rd1, [variadics2_param_1];
; CHECK-PTX-NEXT: add.s64 %rd2, %rd1, 7;
@@ -155,12 +153,12 @@ define dso_local i32 @variadics2(i32 noundef %first, ...) {
; CHECK-PTX-NEXT: ld.local.b32 %r2, [%rd3];
; CHECK-PTX-NEXT: ld.local.s8 %r3, [%rd3+4];
; CHECK-PTX-NEXT: ld.local.b8 %rs1, [%rd3+7];
-; CHECK-PTX-NEXT: st.local.b8 [%SP+2], %rs1;
+; CHECK-PTX-NEXT: st.local.b8 [%SPL+2], %rs1;
; CHECK-PTX-NEXT: ld.local.b8 %rs2, [%rd3+5];
; CHECK-PTX-NEXT: ld.local.b8 %rs3, [%rd3+6];
; CHECK-PTX-NEXT: shl.b16 %rs4, %rs3, 8;
; CHECK-PTX-NEXT: or.b16 %rs5, %rs4, %rs2;
-; CHECK-PTX-NEXT: st.local.b16 [%SP], %rs5;
+; CHECK-PTX-NEXT: st.local.b16 [%SPL], %rs5;
; CHECK-PTX-NEXT: ld.local.b64 %rd4, [%rd3+8];
; CHECK-PTX-NEXT: add.s32 %r4, %r1, %r2;
; CHECK-PTX-NEXT: add.s32 %r5, %r4, %r3;
@@ -208,22 +206,21 @@ define dso_local i32 @bar() {
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot3;
-; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;
; CHECK-PTX-NEXT: ld.global.nc.b8 %rs1, [__const_$_bar_$_s1+7];
-; CHECK-PTX-NEXT: st.local.b8 [%SP+2], %rs1;
+; CHECK-PTX-NEXT: st.local.b8 [%SPL+2], %rs1;
; CHECK-PTX-NEXT: ld.global.nc.b8 %rs2, [__const_$_bar_$_s1+6];
; CHECK-PTX-NEXT: shl.b16 %rs3, %rs2, 8;
; CHECK-PTX-NEXT: ld.global.nc.b8 %rs4, [__const_$_bar_$_s1+5];
; CHECK-PTX-NEXT: or.b16 %rs5, %rs3, %rs4;
-; CHECK-PTX-NEXT: st.local.b16 [%SP], %rs5;
-; CHECK-PTX-NEXT: st.local.b32 [%SP+8], 1;
-; CHECK-PTX-NEXT: st.local.b8 [%SP+12], 1;
-; CHECK-PTX-NEXT: st.local.b64 [%SP+16], 1;
+; CHECK-PTX-NEXT: st.local.b16 [%SPL], %rs5;
+; CHECK-PTX-NEXT: st.local.b32 [%SPL+8], 1;
+; CHECK-PTX-NEXT: st.local.b8 [%SPL+12], 1;
+; CHECK-PTX-NEXT: st.local.b64 [%SPL+16], 1;
; CHECK-PTX-NEXT: { // callseq 1, 0
; CHECK-PTX-NEXT: .param .b32 param0;
; CHECK-PTX-NEXT: .param .b64 param1;
; CHECK-PTX-NEXT: .param .b32 retval0;
-; CHECK-PTX-NEXT: add.u64 %rd1, %SP, 8;
+; CHECK-PTX-NEXT: add.u64 %rd1, %SPL, 8;
; CHECK-PTX-NEXT: st.param.b64 [param1], %rd1;
; CHECK-PTX-NEXT: st.param.b32 [param0], 1;
; CHECK-PTX-NEXT: call.uni (retval0), variadics2, (param0, param1);
@@ -288,13 +285,12 @@ define dso_local i32 @baz() {
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot5;
-; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;
-; CHECK-PTX-NEXT: st.local.v4.b32 [%SP], {1, 1, 1, 1};
+; CHECK-PTX-NEXT: st.local.v4.b32 [%SPL], {1, 1, 1, 1};
; CHECK-PTX-NEXT: { // callseq 2, 0
; CHECK-PTX-NEXT: .param .b32 param0;
; CHECK-PTX-NEXT: .param .b64 param1;
; CHECK-PTX-NEXT: .param .b32 retval0;
-; CHECK-PTX-NEXT: add.u64 %rd1, %SP, 0;
+; CHECK-PTX-NEXT: add.u64 %rd1, %SPL, 0;
; CHECK-PTX-NEXT: st.param.b64 [param1], %rd1;
; CHECK-PTX-NEXT: st.param.b32 [param0], 1;
; CHECK-PTX-NEXT: call.uni (retval0), variadics3, (param0, param1);
@@ -353,21 +349,20 @@ define dso_local void @qux() {
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot7;
-; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;
; CHECK-PTX-NEXT: ld.global.nc.b64 %rd1, [__const_$_qux_$_s+8];
-; CHECK-PTX-NEXT: st.local.b64 [%SP+8], %rd1;
+; CHECK-PTX-NEXT: st.local.b64 [%SPL+8], %rd1;
; CHECK-PTX-NEXT: ld.global.nc.b64 %rd2, [__const_$_qux_$_s];
-; CHECK-PTX-NEXT: st.local.b64 [%SP], %rd2;
-; CHECK-PTX-NEXT: st.local.b64 [%SP+16], 1;
+; CHECK-PTX-NEXT: st.local.b64 [%SPL], %rd2;
+; CHECK-PTX-NEXT: st.local.b64 [%SPL+16], 1;
; CHECK-PTX-NEXT: { // callseq 3, 0
; CHECK-PTX-NEXT: .param .align 8 .b8 param0[16];
; CHECK-PTX-NEXT: .param .b64 param1;
; CHECK-PTX-NEXT: .param .b32 retval0;
-; CHECK-PTX-NEXT: add.u64 %rd3, %SP, 16;
+; CHECK-PTX-NEXT: add.u64 %rd3, %SPL, 16;
; CHECK-PTX-NEXT: st.param.b64 [param1], %rd3;
-; CHECK-PTX-NEXT: ld.b64 %rd4, [%SP+8];
+; CHECK-PTX-NEXT: ld.b64 %rd4, [%SPL+8];
; CHECK-PTX-NEXT: st.param.b64 [param0+8], %rd4;
-; CHECK-PTX-NEXT: ld.b64 %rd5, [%SP];
+; CHECK-PTX-NEXT: ld.b64 %rd5, [%SPL];
; CHECK-PTX-NEXT: st.param.b64 [param0], %rd5;
; CHECK-PTX-NEXT: call.uni (retval0), variadics4, (param0, param1);
; CHECK-PTX-NEXT: } // callseq 3
diff --git a/llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll b/llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll
index 72ff20563a1b3..f879aa7e25fd9 100644
--- a/llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll
+++ b/llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll
@@ -6,11 +6,12 @@
; CHECK: .visible .func use_dbg_declare()
; CHECK: .local .align 8 .b8 __local_depot0[8];
; CHECK: mov.b64 %SPL, __local_depot0;
+; CHECK: add.u64 %rd1, %SPL, 0;
+; CHECK: cvta.local.u64 %rd2, %rd1;
; CHECK: .loc 1 5 3 // t.c:5:3
; CHECK: { // callseq 0, 0
; CHECK: .param .b64 param0;
-; CHECK: add.u64 %rd1, %SP, 0;
-; CHECK: st.param.b64 [param0], %rd1;
+; CHECK: st.param.b64 [param0], %rd2;
; CHECK: call.uni escape_foo, (param0);
; CHECK: } // callseq 0
; CHECK: .loc 1 6 1 // t.c:6:1
>From c972c9ff9db918ee53d588d135f9d23465ca16b2 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Wed, 17 Jun 2026 14:10:02 +0200
Subject: [PATCH 03/15] Format.
---
llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp | 28 +++++++++++-----------
1 file changed, 14 insertions(+), 14 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
index 8483d216489d6..c82e09c604ded 100644
--- a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
@@ -30,9 +30,9 @@
#include "llvm/ADT/SmallVector.h"
#include "llvm/IR/DebugInfo.h"
#include "llvm/IR/Function.h"
+#include "llvm/IR/Instructions.h"
#include "llvm/IR/IntrinsicInst.h"
#include "llvm/IR/Intrinsics.h"
-#include "llvm/IR/Instructions.h"
#include "llvm/IR/Type.h"
#include "llvm/Pass.h"
@@ -83,9 +83,8 @@ static void updateMemIntrinsicDeclaration(MemIntrinsic *MI) {
Tys.push_back(MSI->getLength()->getType());
}
- Function *Decl =
- Intrinsic::getOrInsertDeclaration(MI->getModule(), MI->getIntrinsicID(),
- Tys);
+ Function *Decl = Intrinsic::getOrInsertDeclaration(MI->getModule(),
+ MI->getIntrinsicID(), Tys);
MI->setCalledFunction(Decl);
}
@@ -122,9 +121,9 @@ static void convertPointerUsersToLocal(Value *OldPtr, Value *LocalPtr,
if (auto *GEP = dyn_cast<GetElementPtrInst>(UserInst);
GEP && GEP->getPointerOperand() == OldPtr) {
SmallVector<Value *, 4> Indices(GEP->idx_begin(), GEP->idx_end());
- auto *NewGEP = GetElementPtrInst::Create(
- GEP->getSourceElementType(), LocalPtr, Indices, "",
- GEP->getIterator());
+ auto *NewGEP =
+ GetElementPtrInst::Create(GEP->getSourceElementType(), LocalPtr,
+ Indices, "", GEP->getIterator());
NewGEP->setNoWrapFlags(GEP->getNoWrapFlags());
NewGEP->copyMetadata(*GEP);
NewGEP->setDebugLoc(GEP->getDebugLoc());
@@ -144,15 +143,16 @@ static void convertPointerUsersToLocal(Value *OldPtr, Value *LocalPtr,
if (auto *ASC = dyn_cast<AddrSpaceCastInst>(UserInst);
ASC && ASC->getOperand(0) == OldPtr &&
- ASC->getDestAddressSpace() == ADDRESS_SPACE_LOCAL) {
+ ASC->getDestAddressSpace() == ADDRESS_SPACE_LOCAL) {
ASC->replaceAllUsesWith(LocalPtr);
ASC->eraseFromParent();
continue;
}
if (auto *II = dyn_cast<IntrinsicInst>(UserInst);
- II && (II->getIntrinsicID() == Intrinsic::lifetime_start ||
- II->getIntrinsicID() == Intrinsic::lifetime_end) &&
+ II &&
+ (II->getIntrinsicID() == Intrinsic::lifetime_start ||
+ II->getIntrinsicID() == Intrinsic::lifetime_end) &&
isa<AllocaInst>(LocalPtr)) {
// Lifetime markers must reference an alloca directly, so retarget them
// to the local alloca and update the overloaded declaration.
@@ -226,10 +226,10 @@ bool NVPTXLowerAlloca::runOnFunction(Function &F) {
// For allocas already in ADDRESS_SPACE_LOCAL, we just need
// addrspacecast to ADDRESS_SPACE_GENERIC.
if (AllocAddrSpace == ADDRESS_SPACE_GENERIC) {
- auto *AllocaInLocalAS = new AllocaInst(
- allocaInst->getAllocatedType(), ADDRESS_SPACE_LOCAL,
- allocaInst->getArraySize(), allocaInst->getAlign(), "",
- allocaInst->getIterator());
+ auto *AllocaInLocalAS =
+ new AllocaInst(allocaInst->getAllocatedType(), ADDRESS_SPACE_LOCAL,
+ allocaInst->getArraySize(), allocaInst->getAlign(), "",
+ allocaInst->getIterator());
AllocaInLocalAS->takeName(allocaInst);
AllocaInLocalAS->setDebugLoc(allocaInst->getDebugLoc());
AllocaInLocalAS->copyMetadata(*allocaInst);
>From 8e5737c64528d534ab76761ec2886f0aec90e50a Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Wed, 17 Jun 2026 14:11:26 +0200
Subject: [PATCH 04/15] Update test comments.
---
llvm/test/CodeGen/NVPTX/memcpy-alloca-align.ll | 13 ++++++-------
1 file changed, 6 insertions(+), 7 deletions(-)
diff --git a/llvm/test/CodeGen/NVPTX/memcpy-alloca-align.ll b/llvm/test/CodeGen/NVPTX/memcpy-alloca-align.ll
index 601ce9477bd69..00860a96e5442 100644
--- a/llvm/test/CodeGen/NVPTX/memcpy-alloca-align.ll
+++ b/llvm/test/CodeGen/NVPTX/memcpy-alloca-align.ll
@@ -2,10 +2,9 @@
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_50 | %ptxas-verify %}
; The expansion of small memcpy/memmove/memset can raise the alignment of
-; under-aligned stack objects. Verify this also works on NVPTX, which lowers
-; allocas to the local address space via addrspacecast instructions that hide
-; the frame index, both when the pointer is materialized in the same basic
-; block as the memory operation and when it is defined in a different one.
+; under-aligned stack objects. Verify this also works on NVPTX, where generic
+; allocas are lowered to local-address-space allocas before codegen, both for
+; direct uses and for uses in a different basic block.
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
target triple = "nvptx64-nvidia-cuda"
@@ -54,7 +53,7 @@ exit:
}
; The source object's alignment was raised when expanding the first memcpy;
-; the loads of the second copy can only learn about it through the frame index.
+; the loads of the second copy can only learn about it through the frame object.
define void @memcpy_alloca_raised_src(i1 %c, ptr %p) {
; CHECK-LABEL: memcpy_alloca_raised_src(
; CHECK: .local .align 4 .b8 __local_depot{{[0-9]+}}[16];
@@ -115,8 +114,8 @@ define void @memset_alloca() {
ret void
}
-; InferPtrAlign also looks through the addrspacecast, letting DAGCombiner
-; refine the alignment of under-aligned plain loads and stores.
+; Local-address-space alloca lowering also lets DAGCombiner refine the alignment
+; of under-aligned plain loads and stores through the frame object.
define i32 @underaligned_load_store(i32 %v) {
; CHECK-LABEL: underaligned_load_store(
; CHECK: ld.local.b32
>From 015b00d338db7f1e38cb426d3f98cb410178f0b5 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Thu, 18 Jun 2026 14:25:34 +0200
Subject: [PATCH 05/15] Simplify NVPTXLowerAlloca to alloca + addrspacecast +
RAUW
Replace the manual user rewriting with the minimal lowering: create
a local address space alloca, addrspacecast it back to generic, and
replaceAllUsesWith. NVPTXInferAddressSpaces, which runs right after,
propagates the local address space into the users and folds the cast
away where possible.
Two kinds of users still have to reference the alloca directly and are
retargeted explicitly (a flat loop, no recursion):
- lifetime markers, which the verifier requires to take an alloca
operand;
- debug records, so the variable stays described by its (stable) stack
slot rather than the cvta.local result; replaceAllUsesWith alone would
point them at the cast.
At -O0 InferAddressSpaces does not run, so accesses stay generic through
cvta.local and small copies are no longer chunked; -O>0 is unaffected.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply at anthropic.com>
---
llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp | 275 ++++----------------
llvm/test/CodeGen/NVPTX/lower-alloca.ll | 6 +-
llvm/test/CodeGen/NVPTX/lower-byval-args.ll | 32 +--
llvm/test/CodeGen/NVPTX/vaargs.ll | 242 +++++++++--------
4 files changed, 210 insertions(+), 345 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
index c82e09c604ded..0327e80adf12e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
@@ -6,27 +6,28 @@
//
//===----------------------------------------------------------------------===//
//
-// For generic alloca instructions, create an equivalent alloca in local
-// address space and rewrite uses that can directly use local memory. For
+// Replace each generic alloca with an equivalent alloca in the local address
+// space, followed by an addrspacecast back to generic for its users. For
// example,
//
// %A = alloca i32
-// store i32 0, i32* %A ; emits st.u32
+// store i32 0, ptr %A ; emits st.u32
//
-// will be transformed to
+// is transformed to
//
// %A = alloca i32, addrspace(5)
-// store i32 0, i32 addrspace(5)* %A ; emits st.local.u32
+// %A.generic = addrspacecast ptr addrspace(5) %A to ptr
+// store i32 0, ptr %A.generic
//
-// Uses that require a generic pointer use a single addrspacecast from the
-// local alloca.
+// This gives the alloca a local frame index, which stack lowering addresses
+// through the local frame pointer (%SPL). NVPTXInferAddressSpaces, which runs
+// right after this pass, propagates the local address space into the users and
+// folds the cast away where possible (so the store above becomes st.local.u32).
//
//===----------------------------------------------------------------------===//
#include "MCTargetDesc/NVPTXBaseInfo.h"
#include "NVPTX.h"
-#include "llvm/ADT/DenseMap.h"
-#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/IR/DebugInfo.h"
#include "llvm/IR/Function.h"
@@ -56,229 +57,63 @@ char NVPTXLowerAlloca::ID = 0;
INITIALIZE_PASS(NVPTXLowerAlloca, "nvptx-lower-alloca", "Lower Alloca", false,
false)
-static Value *getOrCreateGenericPtr(Value *LocalPtr,
- DenseMap<Value *, Value *> &GenericPtrs) {
- auto It = GenericPtrs.find(LocalPtr);
- if (It != GenericPtrs.end())
- return It->second;
-
- auto *LocalInst = cast<Instruction>(LocalPtr);
- auto *GenericPtr = new AddrSpaceCastInst(
- LocalPtr, PointerType::get(LocalPtr->getContext(), ADDRESS_SPACE_GENERIC),
- "");
- GenericPtr->insertAfter(LocalInst->getIterator());
- GenericPtrs[LocalPtr] = GenericPtr;
- return GenericPtr;
-}
-
-static void updateMemIntrinsicDeclaration(MemIntrinsic *MI) {
- SmallVector<Type *, 3> Tys;
- if (auto *MTI = dyn_cast<MemTransferInst>(MI)) {
- Tys.push_back(MTI->getRawDest()->getType());
- Tys.push_back(MTI->getRawSource()->getType());
- Tys.push_back(MTI->getLength()->getType());
- } else {
- auto *MSI = cast<MemSetInst>(MI);
- Tys.push_back(MSI->getRawDest()->getType());
- Tys.push_back(MSI->getLength()->getType());
- }
-
- Function *Decl = Intrinsic::getOrInsertDeclaration(MI->getModule(),
- MI->getIntrinsicID(), Tys);
- MI->setCalledFunction(Decl);
-}
-
-static void convertPointerUsersToLocal(Value *OldPtr, Value *LocalPtr,
- DenseMap<Value *, Value *> &GenericPtrs,
- SmallPtrSetImpl<Value *> &Visited) {
- if (!Visited.insert(OldPtr).second)
- return;
-
- // Debug records aren't on the use-list visited below, so retarget them here;
- // otherwise the variable's location is lost when the old alloca is erased.
- SmallVector<DbgVariableRecord *, 2> DbgUsers;
- findDbgUsers(OldPtr, DbgUsers);
- for (DbgVariableRecord *DVR : DbgUsers)
- DVR->replaceVariableLocationOp(OldPtr, LocalPtr);
-
- for (Use &U : llvm::make_early_inc_range(OldPtr->uses())) {
- auto *UserInst = dyn_cast<Instruction>(U.getUser());
- if (!UserInst)
- continue;
-
- if (auto *LI = dyn_cast<LoadInst>(UserInst);
- LI && LI->getPointerOperand() == OldPtr && !LI->isVolatile()) {
- U.set(LocalPtr);
- continue;
- }
-
- if (auto *SI = dyn_cast<StoreInst>(UserInst);
- SI && SI->getPointerOperand() == OldPtr && !SI->isVolatile()) {
- U.set(LocalPtr);
- continue;
- }
-
- if (auto *GEP = dyn_cast<GetElementPtrInst>(UserInst);
- GEP && GEP->getPointerOperand() == OldPtr) {
- SmallVector<Value *, 4> Indices(GEP->idx_begin(), GEP->idx_end());
- auto *NewGEP =
- GetElementPtrInst::Create(GEP->getSourceElementType(), LocalPtr,
- Indices, "", GEP->getIterator());
- NewGEP->setNoWrapFlags(GEP->getNoWrapFlags());
- NewGEP->copyMetadata(*GEP);
- NewGEP->setDebugLoc(GEP->getDebugLoc());
- convertPointerUsersToLocal(GEP, NewGEP, GenericPtrs, Visited);
- if (GEP->use_empty())
- GEP->eraseFromParent();
- continue;
- }
-
- if (auto *BC = dyn_cast<BitCastInst>(UserInst);
- BC && BC->getOperand(0) == OldPtr && BC->getType()->isPointerTy()) {
- convertPointerUsersToLocal(BC, LocalPtr, GenericPtrs, Visited);
- if (BC->use_empty())
- BC->eraseFromParent();
- continue;
- }
-
- if (auto *ASC = dyn_cast<AddrSpaceCastInst>(UserInst);
- ASC && ASC->getOperand(0) == OldPtr &&
- ASC->getDestAddressSpace() == ADDRESS_SPACE_LOCAL) {
- ASC->replaceAllUsesWith(LocalPtr);
- ASC->eraseFromParent();
- continue;
- }
-
- if (auto *II = dyn_cast<IntrinsicInst>(UserInst);
- II &&
- (II->getIntrinsicID() == Intrinsic::lifetime_start ||
- II->getIntrinsicID() == Intrinsic::lifetime_end) &&
- isa<AllocaInst>(LocalPtr)) {
- // Lifetime markers must reference an alloca directly, so retarget them
- // to the local alloca and update the overloaded declaration.
- U.set(LocalPtr);
- Function *Decl = Intrinsic::getOrInsertDeclaration(
- II->getModule(), II->getIntrinsicID(), {LocalPtr->getType()});
- II->setCalledFunction(Decl);
- continue;
- }
-
- if (auto *MI = dyn_cast<MemIntrinsic>(UserInst)) {
- if (auto *MTI = dyn_cast<MemTransferInst>(MI)) {
- bool Changed = false;
- if (MTI->getRawDest() == OldPtr) {
- MTI->getRawDestUse().set(LocalPtr);
- Changed = true;
- }
- if (MTI->getRawSource() == OldPtr) {
- MTI->getRawSourceUse().set(LocalPtr);
- Changed = true;
- }
- if (Changed) {
- updateMemIntrinsicDeclaration(MI);
- continue;
- }
- } else if (auto *MSI = dyn_cast<MemSetInst>(MI);
- MSI && MSI->getRawDest() == OldPtr) {
- MSI->getRawDestUse().set(LocalPtr);
- updateMemIntrinsicDeclaration(MI);
- continue;
- }
- }
-
- U.set(getOrCreateGenericPtr(LocalPtr, GenericPtrs));
- }
-}
-
// =============================================================================
// Main function for this pass.
// =============================================================================
bool NVPTXLowerAlloca::runOnFunction(Function &F) {
// Mandatory lowering: later stack lowering relies on local allocas, so run
// even for optnone functions (skipFunction is intentionally not called).
- bool Changed = false;
SmallVector<AllocaInst *, 8> Allocas;
- for (auto &BB : F) {
- for (auto &I : BB) {
- if (auto *AI = dyn_cast<AllocaInst>(&I))
+ for (auto &BB : F)
+ for (auto &I : BB)
+ if (auto *AI = dyn_cast<AllocaInst>(&I);
+ AI && AI->getAddressSpace() == ADDRESS_SPACE_GENERIC)
Allocas.push_back(AI);
- }
- }
-
- for (AllocaInst *allocaInst : Allocas) {
- Changed = true;
-
- unsigned AllocAddrSpace = allocaInst->getAddressSpace();
- assert((AllocAddrSpace == ADDRESS_SPACE_GENERIC ||
- AllocAddrSpace == ADDRESS_SPACE_LOCAL) &&
- "AllocaInst can only be in Generic or Local address space for "
- "NVPTX.");
- Instruction *AllocaInLocalAS = allocaInst;
-
- // We need to make sure that LLVM has info that alloca needs to go to
- // ADDRESS_SPACE_LOCAL for InferAddressSpace pass.
- //
- // For allocas in ADDRESS_SPACE_GENERIC, create an equivalent local alloca
- // and rewrite localizable users to use it directly. Users that still need a
- // generic pointer use a single addrspacecast from the local alloca.
- //
- // For allocas already in ADDRESS_SPACE_LOCAL, we just need
- // addrspacecast to ADDRESS_SPACE_GENERIC.
- if (AllocAddrSpace == ADDRESS_SPACE_GENERIC) {
- auto *AllocaInLocalAS =
- new AllocaInst(allocaInst->getAllocatedType(), ADDRESS_SPACE_LOCAL,
- allocaInst->getArraySize(), allocaInst->getAlign(), "",
- allocaInst->getIterator());
- AllocaInLocalAS->takeName(allocaInst);
- AllocaInLocalAS->setDebugLoc(allocaInst->getDebugLoc());
- AllocaInLocalAS->copyMetadata(*allocaInst);
- AllocaInLocalAS->setUsedWithInAlloca(allocaInst->isUsedWithInAlloca());
- AllocaInLocalAS->setSwiftError(allocaInst->isSwiftError());
- DenseMap<Value *, Value *> GenericPtrs;
- SmallPtrSet<Value *, 8> Visited;
- convertPointerUsersToLocal(allocaInst, AllocaInLocalAS, GenericPtrs,
- Visited);
- assert(allocaInst->use_empty() &&
- "All generic alloca uses should have been rewritten.");
- allocaInst->eraseFromParent();
- continue;
- }
-
- auto AllocaInGenericAS = new AddrSpaceCastInst(
- AllocaInLocalAS,
- PointerType::get(allocaInst->getContext(), ADDRESS_SPACE_GENERIC), "");
- AllocaInGenericAS->insertAfter(AllocaInLocalAS->getIterator());
-
- for (Use &AllocaUse : llvm::make_early_inc_range(allocaInst->uses())) {
- // Check Load, Store, GEP, and BitCast Uses on alloca and make them
- // use the converted generic address, in order to expose non-generic
- // addrspacecast to NVPTXInferAddressSpaces. For other types
- // of instructions this is unnecessary and may introduce redundant
- // address cast.
- auto LI = dyn_cast<LoadInst>(AllocaUse.getUser());
- if (LI && LI->getPointerOperand() == allocaInst && !LI->isVolatile()) {
- LI->setOperand(LI->getPointerOperandIndex(), AllocaInGenericAS);
- continue;
- }
- auto SI = dyn_cast<StoreInst>(AllocaUse.getUser());
- if (SI && SI->getPointerOperand() == allocaInst && !SI->isVolatile()) {
- SI->setOperand(SI->getPointerOperandIndex(), AllocaInGenericAS);
- continue;
- }
- auto GI = dyn_cast<GetElementPtrInst>(AllocaUse.getUser());
- if (GI && GI->getPointerOperand() == allocaInst) {
- GI->setOperand(GI->getPointerOperandIndex(), AllocaInGenericAS);
+ for (AllocaInst *AI : Allocas) {
+ // Create an equivalent alloca in the local address space.
+ auto *LocalAlloca = new AllocaInst(
+ AI->getAllocatedType(), ADDRESS_SPACE_LOCAL, AI->getArraySize(),
+ AI->getAlign(), "", AI->getIterator());
+ LocalAlloca->setDebugLoc(AI->getDebugLoc());
+ LocalAlloca->copyMetadata(*AI);
+ LocalAlloca->setUsedWithInAlloca(AI->isUsedWithInAlloca());
+ LocalAlloca->setSwiftError(AI->isSwiftError());
+
+ // Debug records and lifetime markers have to reference the alloca itself,
+ // not a cast of it, so retarget them to the local alloca before rewriting
+ // the remaining users through a generic addrspacecast below:
+ // - the verifier requires an alloca operand for lifetime markers, and
+ // - pointing debug records at the alloca keeps the variable described by
+ // its (stable) stack slot rather than the cvta.local result.
+ SmallVector<DbgVariableRecord *, 2> DbgUsers;
+ findDbgUsers(AI, DbgUsers);
+ for (DbgVariableRecord *DVR : DbgUsers)
+ DVR->replaceVariableLocationOp(AI, LocalAlloca);
+
+ for (Use &U : llvm::make_early_inc_range(AI->uses())) {
+ auto *II = dyn_cast<IntrinsicInst>(U.getUser());
+ if (!II || !isLifetimeIntrinsic(II->getIntrinsicID()))
continue;
- }
- auto BI = dyn_cast<BitCastInst>(AllocaUse.getUser());
- if (BI && BI->getOperand(0) == allocaInst) {
- BI->setOperand(0, AllocaInGenericAS);
- continue;
- }
+ U.set(LocalAlloca);
+ Function *Decl = Intrinsic::getOrInsertDeclaration(
+ II->getModule(), II->getIntrinsicID(), {LocalAlloca->getType()});
+ II->setCalledFunction(Decl);
}
+
+ // Everything else can go through a single generic addrspacecast.
+ // replaceAllUsesWith leaves the (already retargeted) lifetime markers and
+ // debug records untouched. NVPTXInferAddressSpaces folds the cast into the
+ // users that can operate on local memory directly.
+ auto *GenericPtr =
+ new AddrSpaceCastInst(LocalAlloca, AI->getType(), "", AI->getIterator());
+ GenericPtr->setDebugLoc(AI->getDebugLoc());
+ AI->replaceAllUsesWith(GenericPtr);
+ LocalAlloca->takeName(AI);
+ AI->eraseFromParent();
}
- return Changed;
+
+ return !Allocas.empty();
}
FunctionPass *llvm::createNVPTXLowerAllocaPass() {
diff --git a/llvm/test/CodeGen/NVPTX/lower-alloca.ll b/llvm/test/CodeGen/NVPTX/lower-alloca.ll
index b93e6a316ddf4..99379d64e5622 100644
--- a/llvm/test/CodeGen/NVPTX/lower-alloca.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-alloca.ll
@@ -16,7 +16,7 @@ define ptx_kernel void @kernel() {
; CHECK: call void @callee(ptr [[GENERIC]])
; LOWERALLOCAONLY: [[A:%.*]] = alloca i32, align 4, addrspace(5)
; LOWERALLOCAONLY: [[GENERIC:%.*]] = addrspacecast ptr addrspace(5) [[A]] to ptr
-; LOWERALLOCAONLY: store i32 0, ptr addrspace(5) [[A]], align 4
+; LOWERALLOCAONLY: store i32 0, ptr [[GENERIC]], align 4
; LOWERALLOCAONLY: call void @callee(ptr [[GENERIC]])
; PTX: st.local.b32 [%SPL], 0
store i32 0, ptr %A
@@ -30,8 +30,8 @@ define void @alloca_in_explicit_local_as() {
%A = alloca i32, addrspace(5)
; CHECK: store i32 0, ptr addrspace(5) {{%.+}}
; PTX: st.local.b32 [%SPL], 0
-; LOWERALLOCAONLY: [[V1:%.*]] = addrspacecast ptr addrspace(5) %A to ptr
-; LOWERALLOCAONLY: store i32 0, ptr [[V1]], align 4
+; An alloca already in the local address space is left as-is.
+; LOWERALLOCAONLY: store i32 0, ptr addrspace(5) %A, align 4
store i32 0, ptr addrspace(5) %A
call void @callee(ptr addrspace(5) %A)
ret void
diff --git a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
index 57c9abe372f48..02a9c4bc29a27 100644
--- a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
@@ -194,14 +194,14 @@ define dso_local ptx_kernel void @escape_ptr_gep(ptr nocapture noundef readnone
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %entry
; PTX-NEXT: mov.b64 %SPL, __local_depot3;
-; PTX-NEXT: ld.param.b32 %rd1, [escape_ptr_gep_param_1+4];
-; PTX-NEXT: shl.b64 %rd2, %rd1, 32;
-; PTX-NEXT: ld.param.b32 %rd3, [escape_ptr_gep_param_1];
-; PTX-NEXT: or.b64 %rd4, %rd2, %rd3;
-; PTX-NEXT: st.local.b64 [%SPL], %rd4;
-; PTX-NEXT: add.u64 %rd5, %SPL, 0;
-; PTX-NEXT: or.b64 %rd6, %rd5, 4;
-; PTX-NEXT: cvta.local.u64 %rd7, %rd6;
+; PTX-NEXT: add.u64 %rd1, %SPL, 0;
+; PTX-NEXT: cvta.local.u64 %rd2, %rd1;
+; PTX-NEXT: ld.param.b32 %rd3, [escape_ptr_gep_param_1+4];
+; PTX-NEXT: shl.b64 %rd4, %rd3, 32;
+; PTX-NEXT: ld.param.b32 %rd5, [escape_ptr_gep_param_1];
+; PTX-NEXT: or.b64 %rd6, %rd4, %rd5;
+; PTX-NEXT: st.local.b64 [%SPL], %rd6;
+; PTX-NEXT: add.s64 %rd7, %rd2, 4;
; PTX-NEXT: { // callseq 1, 0
; PTX-NEXT: .param .b64 param0;
; PTX-NEXT: st.param.b64 [param0], %rd7;
@@ -288,14 +288,14 @@ define dso_local ptx_kernel void @escape_ptr_gep_store(ptr nocapture noundef wri
; PTX-NEXT: mov.b64 %SPL, __local_depot5;
; PTX-NEXT: ld.param.b64 %rd1, [escape_ptr_gep_store_param_0];
; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;
-; PTX-NEXT: ld.param.b32 %rd3, [escape_ptr_gep_store_param_1+4];
-; PTX-NEXT: shl.b64 %rd4, %rd3, 32;
-; PTX-NEXT: ld.param.b32 %rd5, [escape_ptr_gep_store_param_1];
-; PTX-NEXT: or.b64 %rd6, %rd4, %rd5;
-; PTX-NEXT: st.local.b64 [%SPL], %rd6;
-; PTX-NEXT: add.u64 %rd7, %SPL, 0;
-; PTX-NEXT: or.b64 %rd8, %rd7, 4;
-; PTX-NEXT: cvta.local.u64 %rd9, %rd8;
+; PTX-NEXT: add.u64 %rd3, %SPL, 0;
+; PTX-NEXT: cvta.local.u64 %rd4, %rd3;
+; PTX-NEXT: ld.param.b32 %rd5, [escape_ptr_gep_store_param_1+4];
+; PTX-NEXT: shl.b64 %rd6, %rd5, 32;
+; PTX-NEXT: ld.param.b32 %rd7, [escape_ptr_gep_store_param_1];
+; PTX-NEXT: or.b64 %rd8, %rd6, %rd7;
+; PTX-NEXT: st.local.b64 [%SPL], %rd8;
+; PTX-NEXT: add.s64 %rd9, %rd4, 4;
; PTX-NEXT: st.global.b64 [%rd2], %rd9;
; PTX-NEXT: ret;
entry:
diff --git a/llvm/test/CodeGen/NVPTX/vaargs.ll b/llvm/test/CodeGen/NVPTX/vaargs.ll
index 8802ae3be1816..8ff429a16b74c 100644
--- a/llvm/test/CodeGen/NVPTX/vaargs.ll
+++ b/llvm/test/CodeGen/NVPTX/vaargs.ll
@@ -16,41 +16,52 @@ define i32 @foo(i32 %a, ...) {
; CHECK32-NEXT: .local .align 8 .b8 __local_depot0[32];
; CHECK32-NEXT: .reg .b32 %SP;
; CHECK32-NEXT: .reg .b32 %SPL;
-; CHECK32-NEXT: .reg .b32 %r<30>;
+; CHECK32-NEXT: .reg .b16 %rs<5>;
+; CHECK32-NEXT: .reg .b32 %r<33>;
; CHECK32-NEXT: .reg .b64 %rd<3>;
; CHECK32-EMPTY:
; CHECK32-NEXT: // %bb.0: // %entry
; CHECK32-NEXT: mov.b32 %SPL, __local_depot0;
; CHECK32-NEXT: ld.param.b32 %r2, [foo_param_1];
; CHECK32-NEXT: ld.param.b32 %r1, [foo_param_0];
-; CHECK32-NEXT: cvta.local.u32 %r3, %r2;
-; CHECK32-NEXT: st.local.b32 [%SPL], %r3;
-; CHECK32-NEXT: ld.local.b32 %r4, [%SPL];
-; CHECK32-NEXT: st.local.b32 [%SPL+16], %r4;
-; CHECK32-NEXT: ld.local.b32 %r5, [%SPL];
-; CHECK32-NEXT: add.s32 %r6, %r5, 3;
-; CHECK32-NEXT: and.b32 %r7, %r6, -4;
-; CHECK32-NEXT: add.s32 %r8, %r7, 4;
-; CHECK32-NEXT: st.local.b32 [%SPL], %r8;
-; CHECK32-NEXT: ld.b32 %r9, [%r7];
-; CHECK32-NEXT: ld.local.b32 %r10, [%SPL];
-; CHECK32-NEXT: add.s32 %r11, %r10, 7;
-; CHECK32-NEXT: and.b32 %r12, %r11, -8;
-; CHECK32-NEXT: add.s32 %r13, %r12, 8;
-; CHECK32-NEXT: st.local.b32 [%SPL], %r13;
-; CHECK32-NEXT: ld.b64 %rd1, [%r12];
-; CHECK32-NEXT: ld.local.b32 %r14, [%SPL];
-; CHECK32-NEXT: add.s32 %r15, %r14, 7;
-; CHECK32-NEXT: and.b32 %r16, %r15, -8;
-; CHECK32-NEXT: add.s32 %r17, %r16, 8;
-; CHECK32-NEXT: st.local.b32 [%SPL], %r17;
-; CHECK32-NEXT: ld.b64 %rd2, [%r16];
-; CHECK32-NEXT: ld.local.b32 %r18, [%SPL];
-; CHECK32-NEXT: add.s32 %r19, %r18, 3;
-; CHECK32-NEXT: and.b32 %r20, %r19, -4;
-; CHECK32-NEXT: add.s32 %r21, %r20, 4;
-; CHECK32-NEXT: st.local.b32 [%SPL], %r21;
-; CHECK32-NEXT: ld.b32 %r22, [%r20];
+; CHECK32-NEXT: add.u32 %r3, %SPL, 0;
+; CHECK32-NEXT: cvta.local.u32 %r4, %r3;
+; CHECK32-NEXT: add.u32 %r5, %SPL, 16;
+; CHECK32-NEXT: cvta.local.u32 %r6, %r5;
+; CHECK32-NEXT: cvta.local.u32 %r7, %r2;
+; CHECK32-NEXT: st.b32 [%r4], %r7;
+; CHECK32-NEXT: ld.b8 %rs1, [%r4+3];
+; CHECK32-NEXT: st.b8 [%r6+3], %rs1;
+; CHECK32-NEXT: ld.b8 %rs2, [%r4+2];
+; CHECK32-NEXT: st.b8 [%r6+2], %rs2;
+; CHECK32-NEXT: ld.b8 %rs3, [%r4+1];
+; CHECK32-NEXT: st.b8 [%r6+1], %rs3;
+; CHECK32-NEXT: ld.b8 %rs4, [%r4];
+; CHECK32-NEXT: st.b8 [%r6], %rs4;
+; CHECK32-NEXT: ld.b32 %r8, [%r4];
+; CHECK32-NEXT: add.s32 %r9, %r8, 3;
+; CHECK32-NEXT: and.b32 %r10, %r9, -4;
+; CHECK32-NEXT: add.s32 %r11, %r10, 4;
+; CHECK32-NEXT: st.b32 [%r4], %r11;
+; CHECK32-NEXT: ld.b32 %r12, [%r10];
+; CHECK32-NEXT: ld.b32 %r13, [%r4];
+; CHECK32-NEXT: add.s32 %r14, %r13, 7;
+; CHECK32-NEXT: and.b32 %r15, %r14, -8;
+; CHECK32-NEXT: add.s32 %r16, %r15, 8;
+; CHECK32-NEXT: st.b32 [%r4], %r16;
+; CHECK32-NEXT: ld.b64 %rd1, [%r15];
+; CHECK32-NEXT: ld.b32 %r17, [%r4];
+; CHECK32-NEXT: add.s32 %r18, %r17, 7;
+; CHECK32-NEXT: and.b32 %r19, %r18, -8;
+; CHECK32-NEXT: add.s32 %r20, %r19, 8;
+; CHECK32-NEXT: st.b32 [%r4], %r20;
+; CHECK32-NEXT: ld.b64 %rd2, [%r19];
+; CHECK32-NEXT: ld.b32 %r21, [%r4];
+; CHECK32-NEXT: add.s32 %r22, %r21, 3;
+; CHECK32-NEXT: and.b32 %r23, %r22, -4;
+; CHECK32-NEXT: add.s32 %r24, %r23, 4;
+; CHECK32-NEXT: st.b32 [%r4], %r24;
+; CHECK32-NEXT: ld.b32 %r25, [%r23];
; CHECK32-NEXT: { // callseq 0, 0
; CHECK32-NEXT: .param .b32 param0;
; CHECK32-NEXT: .param .b32 param1;
@@ -58,22 +69,22 @@ define i32 @foo(i32 %a, ...) {
; CHECK32-NEXT: .param .b64 param3;
; CHECK32-NEXT: .param .b32 param4;
; CHECK32-NEXT: .param .b32 retval0;
-; CHECK32-NEXT: st.param.b32 [param4], %r22;
+; CHECK32-NEXT: st.param.b32 [param4], %r25;
; CHECK32-NEXT: st.param.b64 [param3], %rd2;
; CHECK32-NEXT: st.param.b64 [param2], %rd1;
-; CHECK32-NEXT: st.param.b32 [param1], %r9;
+; CHECK32-NEXT: st.param.b32 [param1], %r12;
; CHECK32-NEXT: st.param.b32 [param0], %r1;
; CHECK32-NEXT: call.uni (retval0), bar, (param0, param1, param2, param3, param4);
-; CHECK32-NEXT: ld.param.b32 %r23, [retval0];
+; CHECK32-NEXT: ld.param.b32 %r26, [retval0];
; CHECK32-NEXT: } // callseq 0
-; CHECK32-NEXT: ld.local.b32 %r24, [%SPL+16];
-; CHECK32-NEXT: add.s32 %r25, %r24, 3;
-; CHECK32-NEXT: and.b32 %r26, %r25, -4;
-; CHECK32-NEXT: add.s32 %r27, %r26, 4;
-; CHECK32-NEXT: st.local.b32 [%SPL+16], %r27;
-; CHECK32-NEXT: ld.b32 %r28, [%r26];
-; CHECK32-NEXT: add.s32 %r29, %r23, %r28;
-; CHECK32-NEXT: st.param.b32 [func_retval0], %r29;
+; CHECK32-NEXT: ld.b32 %r27, [%r6];
+; CHECK32-NEXT: add.s32 %r28, %r27, 3;
+; CHECK32-NEXT: and.b32 %r29, %r28, -4;
+; CHECK32-NEXT: add.s32 %r30, %r29, 4;
+; CHECK32-NEXT: st.b32 [%r6], %r30;
+; CHECK32-NEXT: ld.b32 %r31, [%r29];
+; CHECK32-NEXT: add.s32 %r32, %r26, %r31;
+; CHECK32-NEXT: st.param.b32 [func_retval0], %r32;
; CHECK32-NEXT: ret;
;
; CHECK64-LABEL: foo(
@@ -81,41 +92,60 @@ define i32 @foo(i32 %a, ...) {
; CHECK64-NEXT: .local .align 8 .b8 __local_depot0[48];
; CHECK64-NEXT: .reg .b64 %SP;
; CHECK64-NEXT: .reg .b64 %SPL;
+; CHECK64-NEXT: .reg .b16 %rs<9>;
; CHECK64-NEXT: .reg .b32 %r<6>;
-; CHECK64-NEXT: .reg .b64 %rd<27>;
+; CHECK64-NEXT: .reg .b64 %rd<30>;
; CHECK64-EMPTY:
; CHECK64-NEXT: // %bb.0: // %entry
; CHECK64-NEXT: mov.b64 %SPL, __local_depot0;
; CHECK64-NEXT: ld.param.b64 %rd1, [foo_param_1];
; CHECK64-NEXT: ld.param.b32 %r1, [foo_param_0];
-; CHECK64-NEXT: cvta.local.u64 %rd2, %rd1;
-; CHECK64-NEXT: st.local.b64 [%SPL], %rd2;
-; CHECK64-NEXT: ld.local.b64 %rd3, [%SPL];
-; CHECK64-NEXT: st.local.b64 [%SPL+24], %rd3;
-; CHECK64-NEXT: ld.local.b64 %rd4, [%SPL];
-; CHECK64-NEXT: add.s64 %rd5, %rd4, 3;
-; CHECK64-NEXT: and.b64 %rd6, %rd5, -4;
-; CHECK64-NEXT: add.s64 %rd7, %rd6, 4;
-; CHECK64-NEXT: st.local.b64 [%SPL], %rd7;
-; CHECK64-NEXT: ld.b32 %r2, [%rd6];
-; CHECK64-NEXT: ld.local.b64 %rd8, [%SPL];
-; CHECK64-NEXT: add.s64 %rd9, %rd8, 7;
-; CHECK64-NEXT: and.b64 %rd10, %rd9, -8;
-; CHECK64-NEXT: add.s64 %rd11, %rd10, 8;
-; CHECK64-NEXT: st.local.b64 [%SPL], %rd11;
-; CHECK64-NEXT: ld.b64 %rd12, [%rd10];
-; CHECK64-NEXT: ld.local.b64 %rd13, [%SPL];
-; CHECK64-NEXT: add.s64 %rd14, %rd13, 7;
-; CHECK64-NEXT: and.b64 %rd15, %rd14, -8;
-; CHECK64-NEXT: add.s64 %rd16, %rd15, 8;
-; CHECK64-NEXT: st.local.b64 [%SPL], %rd16;
-; CHECK64-NEXT: ld.b64 %rd17, [%rd15];
-; CHECK64-NEXT: ld.local.b64 %rd18, [%SPL];
-; CHECK64-NEXT: add.s64 %rd19, %rd18, 7;
-; CHECK64-NEXT: and.b64 %rd20, %rd19, -8;
-; CHECK64-NEXT: add.s64 %rd21, %rd20, 8;
-; CHECK64-NEXT: st.local.b64 [%SPL], %rd21;
-; CHECK64-NEXT: ld.b64 %rd22, [%rd20];
+; CHECK64-NEXT: add.u64 %rd2, %SPL, 0;
+; CHECK64-NEXT: cvta.local.u64 %rd3, %rd2;
+; CHECK64-NEXT: add.u64 %rd4, %SPL, 24;
+; CHECK64-NEXT: cvta.local.u64 %rd5, %rd4;
+; CHECK64-NEXT: cvta.local.u64 %rd6, %rd1;
+; CHECK64-NEXT: st.b64 [%rd3], %rd6;
+; CHECK64-NEXT: ld.b8 %rs1, [%rd3+7];
+; CHECK64-NEXT: st.b8 [%rd5+7], %rs1;
+; CHECK64-NEXT: ld.b8 %rs2, [%rd3+6];
+; CHECK64-NEXT: st.b8 [%rd5+6], %rs2;
+; CHECK64-NEXT: ld.b8 %rs3, [%rd3+5];
+; CHECK64-NEXT: st.b8 [%rd5+5], %rs3;
+; CHECK64-NEXT: ld.b8 %rs4, [%rd3+4];
+; CHECK64-NEXT: st.b8 [%rd5+4], %rs4;
+; CHECK64-NEXT: ld.b8 %rs5, [%rd3+3];
+; CHECK64-NEXT: st.b8 [%rd5+3], %rs5;
+; CHECK64-NEXT: ld.b8 %rs6, [%rd3+2];
+; CHECK64-NEXT: st.b8 [%rd5+2], %rs6;
+; CHECK64-NEXT: ld.b8 %rs7, [%rd3+1];
+; CHECK64-NEXT: st.b8 [%rd5+1], %rs7;
+; CHECK64-NEXT: ld.b8 %rs8, [%rd3];
+; CHECK64-NEXT: st.b8 [%rd5], %rs8;
+; CHECK64-NEXT: ld.b64 %rd7, [%rd3];
+; CHECK64-NEXT: add.s64 %rd8, %rd7, 3;
+; CHECK64-NEXT: and.b64 %rd9, %rd8, -4;
+; CHECK64-NEXT: add.s64 %rd10, %rd9, 4;
+; CHECK64-NEXT: st.b64 [%rd3], %rd10;
+; CHECK64-NEXT: ld.b32 %r2, [%rd9];
+; CHECK64-NEXT: ld.b64 %rd11, [%rd3];
+; CHECK64-NEXT: add.s64 %rd12, %rd11, 7;
+; CHECK64-NEXT: and.b64 %rd13, %rd12, -8;
+; CHECK64-NEXT: add.s64 %rd14, %rd13, 8;
+; CHECK64-NEXT: st.b64 [%rd3], %rd14;
+; CHECK64-NEXT: ld.b64 %rd15, [%rd13];
+; CHECK64-NEXT: ld.b64 %rd16, [%rd3];
+; CHECK64-NEXT: add.s64 %rd17, %rd16, 7;
+; CHECK64-NEXT: and.b64 %rd18, %rd17, -8;
+; CHECK64-NEXT: add.s64 %rd19, %rd18, 8;
+; CHECK64-NEXT: st.b64 [%rd3], %rd19;
+; CHECK64-NEXT: ld.b64 %rd20, [%rd18];
+; CHECK64-NEXT: ld.b64 %rd21, [%rd3];
+; CHECK64-NEXT: add.s64 %rd22, %rd21, 7;
+; CHECK64-NEXT: and.b64 %rd23, %rd22, -8;
+; CHECK64-NEXT: add.s64 %rd24, %rd23, 8;
+; CHECK64-NEXT: st.b64 [%rd3], %rd24;
+; CHECK64-NEXT: ld.b64 %rd25, [%rd23];
; CHECK64-NEXT: { // callseq 0, 0
; CHECK64-NEXT: .param .b32 param0;
; CHECK64-NEXT: .param .b32 param1;
@@ -123,20 +153,20 @@ define i32 @foo(i32 %a, ...) {
; CHECK64-NEXT: .param .b64 param3;
; CHECK64-NEXT: .param .b64 param4;
; CHECK64-NEXT: .param .b32 retval0;
-; CHECK64-NEXT: st.param.b64 [param4], %rd22;
-; CHECK64-NEXT: st.param.b64 [param3], %rd17;
-; CHECK64-NEXT: st.param.b64 [param2], %rd12;
+; CHECK64-NEXT: st.param.b64 [param4], %rd25;
+; CHECK64-NEXT: st.param.b64 [param3], %rd20;
+; CHECK64-NEXT: st.param.b64 [param2], %rd15;
; CHECK64-NEXT: st.param.b32 [param1], %r2;
; CHECK64-NEXT: st.param.b32 [param0], %r1;
; CHECK64-NEXT: call.uni (retval0), bar, (param0, param1, param2, param3, param4);
; CHECK64-NEXT: ld.param.b32 %r3, [retval0];
; CHECK64-NEXT: } // callseq 0
-; CHECK64-NEXT: ld.local.b64 %rd23, [%SPL+24];
-; CHECK64-NEXT: add.s64 %rd24, %rd23, 3;
-; CHECK64-NEXT: and.b64 %rd25, %rd24, -4;
-; CHECK64-NEXT: add.s64 %rd26, %rd25, 4;
-; CHECK64-NEXT: st.local.b64 [%SPL+24], %rd26;
-; CHECK64-NEXT: ld.b32 %r4, [%rd25];
+; CHECK64-NEXT: ld.b64 %rd26, [%rd5];
+; CHECK64-NEXT: add.s64 %rd27, %rd26, 3;
+; CHECK64-NEXT: and.b64 %rd28, %rd27, -4;
+; CHECK64-NEXT: add.s64 %rd29, %rd28, 4;
+; CHECK64-NEXT: st.b64 [%rd5], %rd29;
+; CHECK64-NEXT: ld.b32 %r4, [%rd28];
; CHECK64-NEXT: add.s32 %r5, %r3, %r4;
; CHECK64-NEXT: st.param.b32 [func_retval0], %r5;
; CHECK64-NEXT: ret;
@@ -184,34 +214,34 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK32-NEXT: .local .align 8 .b8 __local_depot1[32];
; CHECK32-NEXT: .reg .b32 %SP;
; CHECK32-NEXT: .reg .b32 %SPL;
-; CHECK32-NEXT: .reg .b32 %r<8>;
+; CHECK32-NEXT: .reg .b32 %r<9>;
; CHECK32-NEXT: .reg .b64 %rd<3>;
-; CHECK32-EMPTY:
+; CHECK32-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _);
; CHECK32-NEXT: // %bb.0: // %entry
; CHECK32-NEXT: mov.b32 %SPL, __local_depot1;
; CHECK32-NEXT: ld.param.b32 %r2, [test_foo_param_3];
; CHECK32-NEXT: ld.param.b64 %rd2, [test_foo_param_2];
; CHECK32-NEXT: ld.param.b64 %rd1, [test_foo_param_1];
; CHECK32-NEXT: ld.param.b32 %r1, [test_foo_param_0];
-; CHECK32-NEXT: mov.b32 %r3, foo_ptr;
-; CHECK32-NEXT: cvta.global.u32 %r4, %r3;
-; CHECK32-NEXT: ld.b32 %r5, [%r4];
-; CHECK32-NEXT: st.local.b32 [%SPL], %r1;
-; CHECK32-NEXT: st.local.b64 [%SPL+8], %rd1;
-; CHECK32-NEXT: st.local.b64 [%SPL+16], %rd2;
-; CHECK32-NEXT: st.local.b32 [%SPL+24], %r2;
+; CHECK32-NEXT: add.u32 %r3, %SPL, 0;
+; CHECK32-NEXT: cvta.local.u32 %r4, %r3;
+; CHECK32-NEXT: mov.b32 %r5, foo_ptr;
+; CHECK32-NEXT: cvta.global.u32 %r6, %r5;
+; CHECK32-NEXT: ld.b32 %r7, [%r6];
+; CHECK32-NEXT: st.b32 [%r4], %r1;
+; CHECK32-NEXT: st.b64 [%r4+8], %rd1;
+; CHECK32-NEXT: st.b64 [%r4+16], %rd2;
+; CHECK32-NEXT: st.b32 [%r4+24], %r2;
; CHECK32-NEXT: { // callseq 1, 0
; CHECK32-NEXT: .param .b32 param0;
; CHECK32-NEXT: .param .b32 param1;
; CHECK32-NEXT: .param .b32 retval0;
-; CHECK32-NEXT: add.u32 %r6, %SPL, 0;
-; CHECK32-NEXT: st.param.b32 [param1], %r6;
-; CHECK32-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _);
+; CHECK32-NEXT: st.param.b32 [param1], %r3;
; CHECK32-NEXT: st.param.b32 [param0], 4;
-; CHECK32-NEXT: call (retval0), %r5, (param0, param1), prototype_1;
-; CHECK32-NEXT: ld.param.b32 %r7, [retval0];
+; CHECK32-NEXT: call (retval0), %r7, (param0, param1), prototype_1;
+; CHECK32-NEXT: ld.param.b32 %r8, [retval0];
; CHECK32-NEXT: } // callseq 1
-; CHECK32-NEXT: st.param.b32 [func_retval0], %r7;
+; CHECK32-NEXT: st.param.b32 [func_retval0], %r8;
; CHECK32-NEXT: ret;
;
; CHECK64-LABEL: test_foo(
@@ -220,30 +250,30 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK64-NEXT: .reg .b64 %SP;
; CHECK64-NEXT: .reg .b64 %SPL;
; CHECK64-NEXT: .reg .b32 %r<3>;
-; CHECK64-NEXT: .reg .b64 %rd<8>;
-; CHECK64-EMPTY:
+; CHECK64-NEXT: .reg .b64 %rd<9>;
+; CHECK64-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b64 _);
; CHECK64-NEXT: // %bb.0: // %entry
; CHECK64-NEXT: mov.b64 %SPL, __local_depot1;
; CHECK64-NEXT: ld.param.b64 %rd3, [test_foo_param_3];
; CHECK64-NEXT: ld.param.b64 %rd2, [test_foo_param_2];
; CHECK64-NEXT: ld.param.b64 %rd1, [test_foo_param_1];
; CHECK64-NEXT: ld.param.b32 %r1, [test_foo_param_0];
-; CHECK64-NEXT: mov.b64 %rd4, foo_ptr;
-; CHECK64-NEXT: cvta.global.u64 %rd5, %rd4;
-; CHECK64-NEXT: ld.b64 %rd6, [%rd5];
-; CHECK64-NEXT: st.local.b32 [%SPL], %r1;
-; CHECK64-NEXT: st.local.b64 [%SPL+8], %rd1;
-; CHECK64-NEXT: st.local.b64 [%SPL+16], %rd2;
-; CHECK64-NEXT: st.local.b64 [%SPL+24], %rd3;
+; CHECK64-NEXT: add.u64 %rd4, %SPL, 0;
+; CHECK64-NEXT: cvta.local.u64 %rd5, %rd4;
+; CHECK64-NEXT: mov.b64 %rd6, foo_ptr;
+; CHECK64-NEXT: cvta.global.u64 %rd7, %rd6;
+; CHECK64-NEXT: ld.b64 %rd8, [%rd7];
+; CHECK64-NEXT: st.b32 [%rd5], %r1;
+; CHECK64-NEXT: st.b64 [%rd5+8], %rd1;
+; CHECK64-NEXT: st.b64 [%rd5+16], %rd2;
+; CHECK64-NEXT: st.b64 [%rd5+24], %rd3;
; CHECK64-NEXT: { // callseq 1, 0
; CHECK64-NEXT: .param .b32 param0;
; CHECK64-NEXT: .param .b64 param1;
; CHECK64-NEXT: .param .b32 retval0;
-; CHECK64-NEXT: add.u64 %rd7, %SPL, 0;
-; CHECK64-NEXT: st.param.b64 [param1], %rd7;
-; CHECK64-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b64 _);
+; CHECK64-NEXT: st.param.b64 [param1], %rd4;
; CHECK64-NEXT: st.param.b32 [param0], 4;
-; CHECK64-NEXT: call (retval0), %rd6, (param0, param1), prototype_1;
+; CHECK64-NEXT: call (retval0), %rd8, (param0, param1), prototype_1;
; CHECK64-NEXT: ld.param.b32 %r2, [retval0];
; CHECK64-NEXT: } // callseq 1
; CHECK64-NEXT: st.param.b32 [func_retval0], %r2;
>From 96cec43dd7ba2d8004c80dad625c3f04580536be Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Thu, 18 Jun 2026 14:42:26 +0200
Subject: [PATCH 06/15] Add lifetime and debug coverage to lower-alloca.ll
The pass-only (no InferAddressSpaces) RUN line mirrors the -O0 pipeline.
Now that the pass defers address-space propagation to IAS, its only
non-trivial logic is retargeting lifetime markers and debug records to the
local alloca. Cover both directly: assert lifetime markers move to the
local alloca with a p5 overload, and that the debug record stays on the
alloca rather than the generic cast.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply at anthropic.com>
---
llvm/test/CodeGen/NVPTX/lower-alloca.ll | 48 +++++++++++++++++++++++++
1 file changed, 48 insertions(+)
diff --git a/llvm/test/CodeGen/NVPTX/lower-alloca.ll b/llvm/test/CodeGen/NVPTX/lower-alloca.ll
index 99379d64e5622..17eff321e9d0e 100644
--- a/llvm/test/CodeGen/NVPTX/lower-alloca.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-alloca.ll
@@ -37,8 +37,56 @@ define void @alloca_in_explicit_local_as() {
ret void
}
+; Lifetime markers must keep referencing the alloca itself (the verifier rejects
+; a cast operand), so they are retargeted to the local alloca and their overload
+; is updated to the local address space, rather than going through the cast.
+define void @lifetime_alloca() {
+; PTX-LABEL: .visible .func lifetime_alloca(
+ %A = alloca i32
+; CHECK: [[A:%.*]] = alloca i32, align 4, addrspace(5)
+; CHECK: call void @llvm.lifetime.start.p5(ptr addrspace(5) [[A]])
+; CHECK: store i32 0, ptr addrspace(5) [[A]]
+; CHECK: call void @llvm.lifetime.end.p5(ptr addrspace(5) [[A]])
+; LOWERALLOCAONLY: [[A:%.*]] = alloca i32, align 4, addrspace(5)
+; LOWERALLOCAONLY: call void @llvm.lifetime.start.p5(ptr addrspace(5) [[A]])
+; LOWERALLOCAONLY: call void @llvm.lifetime.end.p5(ptr addrspace(5) [[A]])
+; PTX: st.local.b32 [%SPL], 0
+ call void @llvm.lifetime.start.p0(ptr %A)
+ store i32 0, ptr %A
+ call void @callee(ptr %A)
+ call void @llvm.lifetime.end.p0(ptr %A)
+ ret void
+}
+
+; Debug records are retargeted to the local alloca too, so the variable stays
+; described by its stack slot rather than the cvta.local result.
+define void @dbg_alloca() !dbg !10 {
+ %A = alloca i32
+; CHECK: [[A:%.*]] = alloca i32, align 4, addrspace(5)
+; CHECK: #dbg_declare(ptr addrspace(5) [[A]],
+; LOWERALLOCAONLY: [[A:%.*]] = alloca i32, align 4, addrspace(5)
+; LOWERALLOCAONLY: #dbg_declare(ptr addrspace(5) [[A]],
+ call void @llvm.dbg.declare(metadata ptr %A, metadata !13, metadata !DIExpression()), !dbg !15
+ store i32 0, ptr %A
+ ret void
+}
+
declare void @callee(ptr)
declare void @callee_addrspace5(ptr addrspace(5))
+declare void @llvm.lifetime.start.p0(ptr)
+declare void @llvm.lifetime.end.p0(ptr)
+declare void @llvm.dbg.declare(metadata, metadata, metadata)
+!llvm.dbg.cu = !{!2}
+!llvm.module.flags = !{!5}
!nvvm.annotations = !{!1}
!1 = !{ptr @alloca_in_explicit_local_as, !"alloca_in_explicit_local_as", i32 1}
+!2 = distinct !DICompileUnit(language: DW_LANG_C99, file: !3, emissionKind: FullDebug)
+!3 = !DIFile(filename: "lower-alloca.c", directory: "/")
+!5 = !{i32 2, !"Debug Info Version", i32 3}
+!10 = distinct !DISubprogram(name: "dbg_alloca", scope: !3, file: !3, line: 1, type: !11, unit: !2)
+!11 = !DISubroutineType(types: !12)
+!12 = !{null}
+!13 = !DILocalVariable(name: "x", scope: !10, file: !3, line: 1, type: !14)
+!14 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
+!15 = !DILocation(line: 1, column: 1, scope: !10)
>From 7275377bc5ff621f8eb68f60f5fc8a7e1dab1458 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Thu, 18 Jun 2026 15:04:26 +0200
Subject: [PATCH 07/15] Clarify lower alloca pass ordering
---
llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
index 0327e80adf12e..003f2e90dc991 100644
--- a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
@@ -20,8 +20,8 @@
// store i32 0, ptr %A.generic
//
// This gives the alloca a local frame index, which stack lowering addresses
-// through the local frame pointer (%SPL). NVPTXInferAddressSpaces, which runs
-// right after this pass, propagates the local address space into the users and
+// through the local frame pointer (%SPL). When NVPTXInferAddressSpaces runs
+// after this pass, it propagates the local address space into the users and
// folds the cast away where possible (so the store above becomes st.local.u32).
//
//===----------------------------------------------------------------------===//
>From 949af1df43f90b0c85a0a1271e5821d51d26498d Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Thu, 18 Jun 2026 15:06:12 +0200
Subject: [PATCH 08/15] Test lower alloca at O0
---
llvm/test/CodeGen/NVPTX/lower-alloca.ll | 5 +++++
1 file changed, 5 insertions(+)
diff --git a/llvm/test/CodeGen/NVPTX/lower-alloca.ll b/llvm/test/CodeGen/NVPTX/lower-alloca.ll
index 17eff321e9d0e..9261c3a99b17d 100644
--- a/llvm/test/CodeGen/NVPTX/lower-alloca.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-alloca.ll
@@ -1,6 +1,7 @@
; RUN: opt < %s -S -nvptx-lower-alloca -infer-address-spaces | FileCheck %s
; RUN: opt < %s -S -nvptx-lower-alloca | FileCheck %s --check-prefix LOWERALLOCAONLY
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_35 | FileCheck %s --check-prefix PTX
+; RUN: llc < %s -O0 -mtriple=nvptx64 -mcpu=sm_35 | FileCheck %s --check-prefix PTXO0
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
@@ -9,6 +10,7 @@ target triple = "nvptx64-unknown-unknown"
define ptx_kernel void @kernel() {
; LABEL: @lower_alloca
; PTX-LABEL: .visible .entry kernel(
+; PTXO0-LABEL: .visible .entry kernel(
%A = alloca i32
; CHECK: [[A:%.*]] = alloca i32, align 4, addrspace(5)
; CHECK: [[GENERIC:%.*]] = addrspacecast ptr addrspace(5) [[A]] to ptr
@@ -19,6 +21,9 @@ define ptx_kernel void @kernel() {
; LOWERALLOCAONLY: store i32 0, ptr [[GENERIC]], align 4
; LOWERALLOCAONLY: call void @callee(ptr [[GENERIC]])
; PTX: st.local.b32 [%SPL], 0
+; PTXO0: mov.b64 %SPL, __local_depot0;
+; PTXO0: cvta.local.u64 [[SP:%rd[0-9]+]], {{%rd[0-9]+}};
+; PTXO0: st.b32 {{\[}}[[SP]]{{\]}}, 0;
store i32 0, ptr %A
call void @callee(ptr %A)
ret void
>From fd2eca8b87bd66586f32b8562503ffc6a753193e Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Thu, 18 Jun 2026 15:32:31 +0200
Subject: [PATCH 09/15] Format.
---
llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp | 10 +++++-----
1 file changed, 5 insertions(+), 5 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
index 003f2e90dc991..760eb401b794d 100644
--- a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
@@ -72,9 +72,9 @@ bool NVPTXLowerAlloca::runOnFunction(Function &F) {
for (AllocaInst *AI : Allocas) {
// Create an equivalent alloca in the local address space.
- auto *LocalAlloca = new AllocaInst(
- AI->getAllocatedType(), ADDRESS_SPACE_LOCAL, AI->getArraySize(),
- AI->getAlign(), "", AI->getIterator());
+ auto *LocalAlloca = new AllocaInst(AI->getAllocatedType(),
+ ADDRESS_SPACE_LOCAL, AI->getArraySize(),
+ AI->getAlign(), "", AI->getIterator());
LocalAlloca->setDebugLoc(AI->getDebugLoc());
LocalAlloca->copyMetadata(*AI);
LocalAlloca->setUsedWithInAlloca(AI->isUsedWithInAlloca());
@@ -105,8 +105,8 @@ bool NVPTXLowerAlloca::runOnFunction(Function &F) {
// replaceAllUsesWith leaves the (already retargeted) lifetime markers and
// debug records untouched. NVPTXInferAddressSpaces folds the cast into the
// users that can operate on local memory directly.
- auto *GenericPtr =
- new AddrSpaceCastInst(LocalAlloca, AI->getType(), "", AI->getIterator());
+ auto *GenericPtr = new AddrSpaceCastInst(LocalAlloca, AI->getType(), "",
+ AI->getIterator());
GenericPtr->setDebugLoc(AI->getDebugLoc());
AI->replaceAllUsesWith(GenericPtr);
LocalAlloca->takeName(AI);
>From e605668f4141ec132e5de4c86cacec588640315d Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Thu, 18 Jun 2026 15:47:31 +0200
Subject: [PATCH 10/15] Shorten comment.
---
llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp | 3 +--
1 file changed, 1 insertion(+), 2 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
index 99b630c02b2c7..7f65eaffb4512 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
@@ -373,8 +373,7 @@ void NVPTXPassConfig::addIRPasses() {
addAddressSpaceInferencePasses();
addStraightLineScalarOptimizationPasses();
} else {
- // Required for correct stack lowering, not just an optimization, so run it
- // even at -O0.
+ // Required for correct stack lowering
addPass(createNVPTXLowerAllocaPass());
}
>From 9bdb7af90b51c41601d5dd77ad42a2f3162090da Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Fri, 10 Jul 2026 10:56:21 +0200
Subject: [PATCH 11/15] Remove dead cast in LowerDYNAMIC_STACKALLOC
NVPTXLowerAlloca guarantees the alloca is local, so the requested value
type is always the local pointer type. Add short-ptr test coverage.
Co-Authored-By: Claude Fable 5 <noreply at anthropic.com>
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 13 +++----
llvm/test/CodeGen/NVPTX/dynamic_stackalloc.ll | 37 +++++++++++++++++++
2 files changed, 42 insertions(+), 8 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index cc7d52b38f5b3..d6c22fcec2ada 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1714,14 +1714,11 @@ SDValue NVPTXTargetLowering::LowerDYNAMIC_STACKALLOC(SDValue Op,
{Chain, DAG.getZExtOrTrunc(Size, DL, LocalVT),
DAG.getTargetConstant(Align, DL, MVT::i32)});
- // Keep the result in the alloca's address space: return the local pointer as
- // is when that is requested (escapes are cast by the caller), else cast it.
- SDValue Ptr = Alloc;
- if (Op.getValueType() != LocalVT)
- Ptr = DAG.getAddrSpaceCast(DL, Op.getValueType(), Alloc,
- ADDRESS_SPACE_LOCAL, ADDRESS_SPACE_GENERIC);
-
- return DAG.getMergeValues({Ptr, SDValue(Alloc.getNode(), 1)}, DL);
+ // NVPTXLowerAlloca puts allocas in the local address space, so a local
+ // pointer is requested here; escapes are explicit addrspacecasts in the IR.
+ assert(Op.getValueType() == LocalVT && "Unexpected alloca pointer size");
+
+ return DAG.getMergeValues({Alloc, SDValue(Alloc.getNode(), 1)}, DL);
}
SDValue NVPTXTargetLowering::LowerSTACKRESTORE(SDValue Op,
diff --git a/llvm/test/CodeGen/NVPTX/dynamic_stackalloc.ll b/llvm/test/CodeGen/NVPTX/dynamic_stackalloc.ll
index 77141277dad2a..4df865585df83 100644
--- a/llvm/test/CodeGen/NVPTX/dynamic_stackalloc.ll
+++ b/llvm/test/CodeGen/NVPTX/dynamic_stackalloc.ll
@@ -4,6 +4,7 @@
; RUN: llc < %s -mtriple=nvptx -mattr=+ptx73 -mcpu=sm_52 | FileCheck %s --check-prefixes=CHECK-32
; RUN: llc < %s -mtriple=nvptx64 -mattr=+ptx73 -mcpu=sm_52 | FileCheck %s --check-prefixes=CHECK-64
+; RUN: llc < %s -mtriple=nvptx64 -mattr=+ptx73 -mcpu=sm_52 --nvptx-short-ptr | FileCheck %s --check-prefixes=CHECK-MIXED
; RUN: %if ptxas-isa-7.3 && ptxas-ptr32 %{ llc < %s -mtriple=nvptx -mattr=+ptx73 -mcpu=sm_52 | %ptxas-verify %}
; RUN: %if ptxas-isa-7.3 %{ llc < %s -mtriple=nvptx64 -mattr=+ptx73 -mcpu=sm_52 | %ptxas-verify %}
@@ -50,6 +51,28 @@ define i32 @test_dynamic_stackalloc(i64 %n) {
; CHECK-64-NEXT: } // callseq 0
; CHECK-64-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-64-NEXT: ret;
+;
+; CHECK-MIXED-LABEL: test_dynamic_stackalloc(
+; CHECK-MIXED: {
+; CHECK-MIXED-NEXT: .reg .b32 %r<6>;
+; CHECK-MIXED-NEXT: .reg .b64 %rd<3>;
+; CHECK-MIXED-EMPTY:
+; CHECK-MIXED-NEXT: // %bb.0:
+; CHECK-MIXED-NEXT: ld.param.b32 %r1, [test_dynamic_stackalloc_param_0];
+; CHECK-MIXED-NEXT: add.s32 %r2, %r1, 7;
+; CHECK-MIXED-NEXT: and.b32 %r3, %r2, -8;
+; CHECK-MIXED-NEXT: alloca.u32 %r4, %r3, 16;
+; CHECK-MIXED-NEXT: cvt.u64.u32 %rd1, %r4;
+; CHECK-MIXED-NEXT: cvta.local.u64 %rd2, %rd1;
+; CHECK-MIXED-NEXT: { // callseq 0, 0
+; CHECK-MIXED-NEXT: .param .b64 param0;
+; CHECK-MIXED-NEXT: .param .b32 retval0;
+; CHECK-MIXED-NEXT: st.param.b64 [param0], %rd2;
+; CHECK-MIXED-NEXT: call.uni (retval0), bar, (param0);
+; CHECK-MIXED-NEXT: ld.param.b32 %r5, [retval0];
+; CHECK-MIXED-NEXT: } // callseq 0
+; CHECK-MIXED-NEXT: st.param.b32 [func_retval0], %r5;
+; CHECK-MIXED-NEXT: ret;
%alloca = alloca i8, i64 %n, align 16
%call = call i32 @bar(ptr %alloca)
ret i32 %call
@@ -84,6 +107,20 @@ define float @test_dynamic_stackalloc_unaligned(i64 %0) {
; CHECK-64-NEXT: ld.local.b32 %r1, [%rd5];
; CHECK-64-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-64-NEXT: ret;
+;
+; CHECK-MIXED-LABEL: test_dynamic_stackalloc_unaligned(
+; CHECK-MIXED: {
+; CHECK-MIXED-NEXT: .reg .b32 %r<7>;
+; CHECK-MIXED-EMPTY:
+; CHECK-MIXED-NEXT: // %bb.0:
+; CHECK-MIXED-NEXT: ld.param.b32 %r1, [test_dynamic_stackalloc_unaligned_param_0];
+; CHECK-MIXED-NEXT: shl.b32 %r2, %r1, 2;
+; CHECK-MIXED-NEXT: add.s32 %r3, %r2, 7;
+; CHECK-MIXED-NEXT: and.b32 %r4, %r3, -8;
+; CHECK-MIXED-NEXT: alloca.u32 %r5, %r4, 8;
+; CHECK-MIXED-NEXT: ld.local.b32 %r6, [%r5];
+; CHECK-MIXED-NEXT: st.param.b32 [func_retval0], %r6;
+; CHECK-MIXED-NEXT: ret;
%4 = alloca float, i64 %0, align 4
%5 = getelementptr float, ptr %4, i64 0
%6 = load float, ptr %5, align 4
>From 7c5d6ece593ee3e639c22e7c49c7ffda505177e1 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Fri, 17 Jul 2026 15:33:53 +0200
Subject: [PATCH 12/15] [NVPTX] Preserve address space for byval copy loads
Use the underlying alloca as pointer information for frame-index loads so SelectionDAG does not replace AS5 with generic fixed-stack information.
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 33 ++++++++++----------
llvm/test/CodeGen/NVPTX/indirect_byval.ll | 4 +--
llvm/test/CodeGen/NVPTX/lower-byval-args.ll | 5 ++-
llvm/test/CodeGen/NVPTX/variadics-backend.ll | 8 ++---
4 files changed, 23 insertions(+), 27 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index d6c22fcec2ada..ee6fcc758e370 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -29,6 +29,7 @@
#include "llvm/ADT/StringRef.h"
#include "llvm/CodeGen/Analysis.h"
#include "llvm/CodeGen/ISDOpcodes.h"
+#include "llvm/CodeGen/MachineFrameInfo.h"
#include "llvm/CodeGen/MachineFunction.h"
#include "llvm/CodeGen/MachineJumpTableInfo.h"
#include "llvm/CodeGen/MachineMemOperand.h"
@@ -1231,28 +1232,25 @@ SDValue NVPTXTargetLowering::getSqrtEstimate(SDValue Operand, SelectionDAG &DAG,
}
}
-static MachinePointerInfo refinePtrAS(SDValue &Ptr, SelectionDAG &DAG,
- const DataLayout &DL,
- const TargetLowering &TL) {
- if (Ptr->getOpcode() == ISD::FrameIndex) {
- auto Ty = TL.getPointerTy(DL, ADDRESS_SPACE_LOCAL);
- Ptr = DAG.getAddrSpaceCast(SDLoc(), Ty, Ptr, ADDRESS_SPACE_GENERIC,
- ADDRESS_SPACE_LOCAL);
-
- return MachinePointerInfo(ADDRESS_SPACE_LOCAL);
- }
-
- // Peel of an addrspacecast to generic and load directly from the specific
- // address space.
+static MachinePointerInfo refinePtrAS(SDValue &Ptr, SelectionDAG &DAG) {
+ // Load directly from the source address space of a cast to generic.
+ unsigned SrcAS = ADDRESS_SPACE_GENERIC;
if (Ptr->getOpcode() == ISD::ADDRSPACECAST) {
const auto *ASC = cast<AddrSpaceCastSDNode>(Ptr);
if (ASC->getDestAddressSpace() == ADDRESS_SPACE_GENERIC) {
Ptr = ASC->getOperand(0);
- return MachinePointerInfo(ASC->getSrcAddressSpace());
+ SrcAS = ASC->getSrcAddressSpace();
}
}
- return MachinePointerInfo();
+ // Preserve the alloca's address space through frame-index inference.
+ if (const auto *FIN = dyn_cast<FrameIndexSDNode>(Ptr))
+ if (const AllocaInst *AI =
+ DAG.getMachineFunction().getFrameInfo().getObjectAllocation(
+ FIN->getIndex()))
+ return MachinePointerInfo(AI);
+
+ return MachinePointerInfo(SrcAS);
}
static ISD::NodeType getExtOpcode(const ISD::ArgFlagsTy &Flags) {
@@ -1419,7 +1417,7 @@ SDValue NVPTXTargetLowering::LowerCall(TargetLowering::CallLoweringInfo &CLI,
if (IsByVal) {
assert(ArgOutVals.size() == 1 && "We must pass only one value as byval");
SDValue SrcPtr = ArgOutVals[0];
- const auto PointerInfo = refinePtrAS(SrcPtr, DAG, DL, *this);
+ const MachinePointerInfo SrcPtrInfo = refinePtrAS(SrcPtr, DAG);
// Don't use Flags.getNonZeroByValAlign as this includes the stackalign,
// which does not apply to the source pointer.
const Align BaseSrcAlign = [&]() {
@@ -1448,7 +1446,8 @@ SDValue NVPTXTargetLowering::LowerCall(TargetLowering::CallLoweringInfo &CLI,
Align SrcAlign = commonAlignment(BaseSrcAlign, Offsets[J]);
SDValue SrcAddr = DAG.getObjectPtrOffset(dl, SrcPtr, Offsets[J]);
SDValue SrcLoad =
- DAG.getLoad(LoadVT, dl, CallChain, SrcAddr, PointerInfo, SrcAlign);
+ DAG.getLoad(LoadVT, dl, CallChain, SrcAddr,
+ SrcPtrInfo.getWithOffset(Offsets[J]), SrcAlign);
TypeSize ParamOffset = Offsets[J].getWithIncrement(VAOffset);
Align ParamAlign = commonAlignment(ArgAlign, ParamOffset);
diff --git a/llvm/test/CodeGen/NVPTX/indirect_byval.ll b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
index 9b8f1f2f8346f..2407c85f3e47e 100644
--- a/llvm/test/CodeGen/NVPTX/indirect_byval.ll
+++ b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
@@ -29,7 +29,7 @@ define internal i32 @foo() {
; CHECK-NEXT: .param .b64 param1;
; CHECK-NEXT: .param .b32 retval0;
; CHECK-NEXT: st.param.b64 [param1], %rd2;
-; CHECK-NEXT: ld.b8 %rs1, [%SPL+1];
+; CHECK-NEXT: ld.local.b8 %rs1, [%SPL+1];
; CHECK-NEXT: st.param.b8 [param0], %rs1;
; CHECK-NEXT: call (retval0), %rd3, (param0, param1), prototype_0;
; CHECK-NEXT: ld.param.b32 %r1, [retval0];
@@ -64,7 +64,7 @@ define internal i32 @bar() {
; CHECK-NEXT: .param .b64 param1;
; CHECK-NEXT: .param .b32 retval0;
; CHECK-NEXT: st.param.b64 [param1], %rd2;
-; CHECK-NEXT: ld.b64 %rd4, [%SPL+8];
+; CHECK-NEXT: ld.local.b64 %rd4, [%SPL+8];
; CHECK-NEXT: st.param.b64 [param0], %rd4;
; CHECK-NEXT: call (retval0), %rd3, (param0, param1), prototype_1;
; CHECK-NEXT: ld.param.b32 %r1, [retval0];
diff --git a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
index 02a9c4bc29a27..ecfa50a92bd13 100644
--- a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
@@ -896,7 +896,7 @@ define ptx_kernel void @test_forward_byval_arg(ptr byval(i32) align 4 %input) {
; PTX-NEXT: .local .align 4 .b8 __local_depot15[4];
; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
-; PTX-NEXT: .reg .b32 %r<3>;
+; PTX-NEXT: .reg .b32 %r<2>;
; PTX-EMPTY:
; PTX-NEXT: // %bb.0:
; PTX-NEXT: mov.b64 %SPL, __local_depot15;
@@ -904,8 +904,7 @@ define ptx_kernel void @test_forward_byval_arg(ptr byval(i32) align 4 %input) {
; PTX-NEXT: st.local.b32 [%SPL], %r1;
; PTX-NEXT: { // callseq 2, 0
; PTX-NEXT: .param .align 4 .b8 param0[4];
-; PTX-NEXT: ld.b32 %r2, [%SPL];
-; PTX-NEXT: st.param.b32 [param0], %r2;
+; PTX-NEXT: st.param.b32 [param0], %r1;
; PTX-NEXT: call.uni device_func, (param0);
; PTX-NEXT: } // callseq 2
; PTX-NEXT: ret;
diff --git a/llvm/test/CodeGen/NVPTX/variadics-backend.ll b/llvm/test/CodeGen/NVPTX/variadics-backend.ll
index e010e71459bca..08f316a22b48c 100644
--- a/llvm/test/CodeGen/NVPTX/variadics-backend.ll
+++ b/llvm/test/CodeGen/NVPTX/variadics-backend.ll
@@ -345,7 +345,7 @@ define dso_local void @qux() {
; CHECK-PTX-NEXT: .local .align 8 .b8 __local_depot7[24];
; CHECK-PTX-NEXT: .reg .b64 %SP;
; CHECK-PTX-NEXT: .reg .b64 %SPL;
-; CHECK-PTX-NEXT: .reg .b64 %rd<6>;
+; CHECK-PTX-NEXT: .reg .b64 %rd<4>;
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot7;
@@ -360,10 +360,8 @@ define dso_local void @qux() {
; CHECK-PTX-NEXT: .param .b32 retval0;
; CHECK-PTX-NEXT: add.u64 %rd3, %SPL, 16;
; CHECK-PTX-NEXT: st.param.b64 [param1], %rd3;
-; CHECK-PTX-NEXT: ld.b64 %rd4, [%SPL+8];
-; CHECK-PTX-NEXT: st.param.b64 [param0+8], %rd4;
-; CHECK-PTX-NEXT: ld.b64 %rd5, [%SPL];
-; CHECK-PTX-NEXT: st.param.b64 [param0], %rd5;
+; CHECK-PTX-NEXT: st.param.b64 [param0+8], %rd1;
+; CHECK-PTX-NEXT: st.param.b64 [param0], %rd2;
; CHECK-PTX-NEXT: call.uni (retval0), variadics4, (param0, param1);
; CHECK-PTX-NEXT: } // callseq 3
; CHECK-PTX-NEXT: ret;
>From 701ab2e57716373f925e4d51fe8f127ee9e5f0a6 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Fri, 17 Jul 2026 15:38:00 +0200
Subject: [PATCH 13/15] [NVPTX] Clarify generic alloca collection
---
llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp | 8 ++++----
1 file changed, 4 insertions(+), 4 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
index 760eb401b794d..38ac883e4980a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
@@ -63,14 +63,14 @@ INITIALIZE_PASS(NVPTXLowerAlloca, "nvptx-lower-alloca", "Lower Alloca", false,
bool NVPTXLowerAlloca::runOnFunction(Function &F) {
// Mandatory lowering: later stack lowering relies on local allocas, so run
// even for optnone functions (skipFunction is intentionally not called).
- SmallVector<AllocaInst *, 8> Allocas;
+ SmallVector<AllocaInst *, 8> GenericAllocas;
for (auto &BB : F)
for (auto &I : BB)
if (auto *AI = dyn_cast<AllocaInst>(&I);
AI && AI->getAddressSpace() == ADDRESS_SPACE_GENERIC)
- Allocas.push_back(AI);
+ GenericAllocas.push_back(AI);
- for (AllocaInst *AI : Allocas) {
+ for (AllocaInst *AI : GenericAllocas) {
// Create an equivalent alloca in the local address space.
auto *LocalAlloca = new AllocaInst(AI->getAllocatedType(),
ADDRESS_SPACE_LOCAL, AI->getArraySize(),
@@ -113,7 +113,7 @@ bool NVPTXLowerAlloca::runOnFunction(Function &F) {
AI->eraseFromParent();
}
- return !Allocas.empty();
+ return !GenericAllocas.empty();
}
FunctionPass *llvm::createNVPTXLowerAllocaPass() {
>From ce56fecea7aa7cfb3fccbdd57431a8ae3bf760eb Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Fri, 17 Jul 2026 15:38:53 +0200
Subject: [PATCH 14/15] [NVPTX] Make indirect byval optimization explicit
llc defaults to -O2, but spelling it out makes the address-space inference coverage clear.
---
llvm/test/CodeGen/NVPTX/indirect_byval.ll | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/llvm/test/CodeGen/NVPTX/indirect_byval.ll b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
index 2407c85f3e47e..7a13bcd1f6674 100644
--- a/llvm/test/CodeGen/NVPTX/indirect_byval.ll
+++ b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
@@ -1,6 +1,6 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
-; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_52 -mattr=+ptx64 | FileCheck %s
-; RUN: %if ptxas-isa-6.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_52 -mattr=+ptx64 | %ptxas-verify %}
+; RUN: llc -O2 < %s -mtriple=nvptx64 -mcpu=sm_52 -mattr=+ptx64 | FileCheck %s
+; RUN: %if ptxas-isa-6.4 %{ llc -O2 < %s -mtriple=nvptx64 -mcpu=sm_52 -mattr=+ptx64 | %ptxas-verify %}
target triple = "nvptx64-nvidia-cuda"
>From c9039cfe7b3fb564ceaca176c086655cea833261 Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Wed, 22 Jul 2026 23:19:11 +0200
Subject: [PATCH 15/15] [NVPTX] Update gridconstant checks for local allocas
---
.../CodeGen/NVPTX/lower-args-gridconstant.ll | 20 +++++++++----------
1 file changed, 9 insertions(+), 11 deletions(-)
diff --git a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
index 33264a1447872..cd51886a3d757 100644
--- a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
@@ -204,22 +204,21 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
; PTX-DEFAULT-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
; PTX-DEFAULT-NEXT: // %bb.0:
; PTX-DEFAULT-NEXT: mov.b64 %SPL, __local_depot4;
-; PTX-DEFAULT-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-DEFAULT-NEXT: mov.b64 %rd1, multiple_grid_const_escape_param_0;
; PTX-DEFAULT-NEXT: mov.b64 %rd2, multiple_grid_const_escape_param_2;
; PTX-DEFAULT-NEXT: cvta.param.u64 %rd3, %rd2;
; PTX-DEFAULT-NEXT: ld.param.b32 %r1, [multiple_grid_const_escape_param_1];
; PTX-DEFAULT-NEXT: cvta.param.u64 %rd4, %rd1;
-; PTX-DEFAULT-NEXT: add.u64 %rd5, %SP, 0;
-; PTX-DEFAULT-NEXT: add.u64 %rd6, %SPL, 0;
-; PTX-DEFAULT-NEXT: st.local.b32 [%rd6], %r1;
+; PTX-DEFAULT-NEXT: add.u64 %rd5, %SPL, 0;
+; PTX-DEFAULT-NEXT: cvta.local.u64 %rd6, %rd5;
+; PTX-DEFAULT-NEXT: st.local.b32 [%SPL], %r1;
; PTX-DEFAULT-NEXT: { // callseq 1, 0
; PTX-DEFAULT-NEXT: .param .b64 param0;
; PTX-DEFAULT-NEXT: .param .b64 param1;
; PTX-DEFAULT-NEXT: .param .b64 param2;
; PTX-DEFAULT-NEXT: .param .b32 retval0;
; PTX-DEFAULT-NEXT: st.param.b64 [param2], %rd3;
-; PTX-DEFAULT-NEXT: st.param.b64 [param1], %rd5;
+; PTX-DEFAULT-NEXT: st.param.b64 [param1], %rd6;
; PTX-DEFAULT-NEXT: st.param.b64 [param0], %rd4;
; PTX-DEFAULT-NEXT: mov.b64 %rd7, escape3;
; PTX-DEFAULT-NEXT: call (retval0), %rd7, (param0, param1, param2), prototype_1;
@@ -236,7 +235,6 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
; PTX-SHORT-PTR-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
; PTX-SHORT-PTR-NEXT: // %bb.0:
; PTX-SHORT-PTR-NEXT: mov.b64 %SPL, __local_depot4;
-; PTX-SHORT-PTR-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, multiple_grid_const_escape_param_0;
; PTX-SHORT-PTR-NEXT: mov.b32 %r2, multiple_grid_const_escape_param_2;
; PTX-SHORT-PTR-NEXT: cvt.u64.u32 %rd1, %r2;
@@ -244,17 +242,17 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
; PTX-SHORT-PTR-NEXT: ld.param.b32 %r3, [multiple_grid_const_escape_param_1];
; PTX-SHORT-PTR-NEXT: cvt.u64.u32 %rd3, %r1;
; PTX-SHORT-PTR-NEXT: cvta.param.u64 %rd4, %rd3;
-; PTX-SHORT-PTR-NEXT: add.u64 %rd5, %SP, 0;
-; PTX-SHORT-PTR-NEXT: add.u64 %rd6, %SPL, 0;
-; PTX-SHORT-PTR-NEXT: cvt.u32.u64 %r4, %rd6;
-; PTX-SHORT-PTR-NEXT: st.local.b32 [%r4], %r3;
+; PTX-SHORT-PTR-NEXT: add.u32 %r4, %SPL, 0;
+; PTX-SHORT-PTR-NEXT: cvt.u64.u32 %rd5, %r4;
+; PTX-SHORT-PTR-NEXT: cvta.local.u64 %rd6, %rd5;
+; PTX-SHORT-PTR-NEXT: st.local.b32 [%SPL], %r3;
; PTX-SHORT-PTR-NEXT: { // callseq 1, 0
; PTX-SHORT-PTR-NEXT: .param .b64 param0;
; PTX-SHORT-PTR-NEXT: .param .b64 param1;
; PTX-SHORT-PTR-NEXT: .param .b64 param2;
; PTX-SHORT-PTR-NEXT: .param .b32 retval0;
; PTX-SHORT-PTR-NEXT: st.param.b64 [param2], %rd2;
-; PTX-SHORT-PTR-NEXT: st.param.b64 [param1], %rd5;
+; PTX-SHORT-PTR-NEXT: st.param.b64 [param1], %rd6;
; PTX-SHORT-PTR-NEXT: st.param.b64 [param0], %rd4;
; PTX-SHORT-PTR-NEXT: mov.b64 %rd7, escape3;
; PTX-SHORT-PTR-NEXT: call (retval0), %rd7, (param0, param1, param2), prototype_1;
More information about the llvm-commits
mailing list