[llvm] [NVPTX] Lower allocas to the local address space (PR #204346)
Tim Besard via llvm-commits
llvm-commits at lists.llvm.org
Wed Jul 8 06:14:46 PDT 2026
https://github.com/maleadt updated https://github.com/llvm/llvm-project/pull/204346
>From 13c6d85e4aaf0f1cbbebc30d6ad4cf4edc2d1c2c 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/10] 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 | 138 +++++----
.../test/CodeGen/NVPTX/memcpy-alloca-align.ll | 132 +++++++++
llvm/test/CodeGen/NVPTX/variadics-backend.ll | 107 +++----
5 files changed, 467 insertions(+), 200 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 8c92264959583..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:
@@ -849,7 +848,7 @@ define ptx_kernel void @test_phi_write(ptr byval(%struct.S) align 4 %input1, ptr
; PTX-NEXT: .reg .pred %p<2>;
; PTX-NEXT: .reg .b16 %rs<3>;
; PTX-NEXT: .reg .b32 %r<3>;
-; PTX-NEXT: .reg .b64 %rd<3>;
+; PTX-NEXT: .reg .b64 %rd<2>;
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %bb
; PTX-NEXT: mov.b64 %SPL, __local_depot14;
@@ -857,17 +856,16 @@ define ptx_kernel void @test_phi_write(ptr byval(%struct.S) align 4 %input1, ptr
; 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: add.u64 %rd1, %SPL, 0;
; PTX-NEXT: ld.param.b32 %r1, [test_phi_write_param_1+4];
-; PTX-NEXT: st.b32 [%SP], %r1;
-; PTX-NEXT: add.u64 %rd2, %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: mov.b64 %rd2, %rd1;
+; PTX-NEXT: add.u64 %rd1, %SP, 0;
; PTX-NEXT: $L__BB14_2: // %merge
-; PTX-NEXT: st.local.b32 [%rd2], 1;
+; PTX-NEXT: st.local.b32 [%rd1], 1;
; PTX-NEXT: ret;
bb:
br i1 %cond, label %first, label %second
@@ -906,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 4b11dd04200fe69085c81bfa10ed94cbc0f18834 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/10] 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 +-
.../CodeGen/NVPTX/lower-args-gridconstant.ll | 9 +-
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 +-
15 files changed, 203 insertions(+), 193 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index fa0b9fe5f7051..efa26847ee023 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1686,10 +1686,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 2bc2158fccae1..438c79208ef7d 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-args-gridconstant.ll b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
index da8c440512af3..d5ad54a3d4372 100644
--- a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
@@ -155,22 +155,21 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
; PTX-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
; PTX-NEXT: // %bb.0:
; PTX-NEXT: mov.b64 %SPL, __local_depot4;
-; PTX-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-NEXT: mov.b64 %rd1, multiple_grid_const_escape_param_0;
; PTX-NEXT: mov.b64 %rd2, multiple_grid_const_escape_param_2;
; PTX-NEXT: cvta.param.u64 %rd3, %rd2;
; PTX-NEXT: ld.param.b32 %r1, [multiple_grid_const_escape_param_1];
; PTX-NEXT: cvta.param.u64 %rd4, %rd1;
-; PTX-NEXT: add.u64 %rd5, %SP, 0;
-; PTX-NEXT: add.u64 %rd6, %SPL, 0;
-; PTX-NEXT: st.local.b32 [%rd6], %r1;
+; PTX-NEXT: add.u64 %rd5, %SPL, 0;
+; PTX-NEXT: cvta.local.u64 %rd6, %rd5;
+; PTX-NEXT: st.local.b32 [%SPL], %r1;
; PTX-NEXT: { // callseq 1, 0
; PTX-NEXT: .param .b64 param0;
; PTX-NEXT: .param .b64 param1;
; PTX-NEXT: .param .b64 param2;
; PTX-NEXT: .param .b32 retval0;
; PTX-NEXT: st.param.b64 [param2], %rd3;
-; PTX-NEXT: st.param.b64 [param1], %rd5;
+; PTX-NEXT: st.param.b64 [param1], %rd6;
; PTX-NEXT: st.param.b64 [param0], %rd4;
; PTX-NEXT: mov.b64 %rd7, escape3;
; PTX-NEXT: call (retval0), %rd7, (param0, param1, param2), prototype_1;
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 af344b7a7b870b24f5692ffc83f39da285a120a7 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/10] 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 4f4f88b21529497a0c9df1283bee3d6db9e9a6aa 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/10] 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 8d64435e24852a1aa5eb7cb15a36d90ff68300c5 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/10] 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 64a14500effeea2f6b5f320dc40c5c796e27e647 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/10] 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 faf47234366ba005d35974ce7c34b69855164f99 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/10] 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 b490121dfabcee601efcaf0b0b7f660132695abd 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/10] 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 b84c83c472b6e1160a547321b89a48877e952668 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/10] 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 c43f37a003d105bd0502ae5456c455cd87b74456 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/10] 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 438c79208ef7d..21d4f7e178818 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());
}
More information about the llvm-commits
mailing list