https://github.com/thetheodor updated
https://github.com/llvm/llvm-project/pull/154814
>From 095073e292e212b938ae3e72c403faab09940a9c Mon Sep 17 00:00:00 2001
From: Theodoros Theodoridis <ttheodori...@nvidia.com>
Date: Tue, 12 Aug 2025 15:22:24 +0000
Subject: [PATCH 1/7] [NVPTX] Change the alloca address space in
NVPTXLowerAlloca
This patch refactors NVPTXLowerAlloca to produce simpler IR for allocas.
Previously, the implementation attached a pair of consecutive address
space casts to each alloca: one from addrspace(0) (generic) to
addrspace(5) (local), and another immediately back to addrspace(0).
Downstream passes needed to recognize this idiom to generate efficient
PTX. With this patch, NVPTXLowerAlloca directly changes the address
space of each alloca to "local" and inserts a single addrspacecast from
local back to generic. The InferAddressSpace pass can then remove the
remaining cast. This change results in fewer address-space-change (ctva)
instructions in the final PTX.
---
llvm/include/llvm/Target/TargetMachine.h | 2 +-
llvm/lib/IR/Verifier.cpp | 8 +
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 3 +-
llvm/lib/Target/NVPTX/NVPTXFrameLowering.cpp | 4 +-
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 5 +-
llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp | 124 ++++------
llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp | 3 +-
llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp | 10 +-
llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp | 24 ++
llvm/lib/Target/NVPTX/NVPTXTargetMachine.h | 3 +
.../CodeGen/NVPTX/call-with-alloca-buffer.ll | 5 +-
.../NVPTX/dynamic-stackalloc-regression.ll | 14 +-
llvm/test/CodeGen/NVPTX/f32x2-instructions.ll | 7 +-
llvm/test/CodeGen/NVPTX/indirect_byval.ll | 22 +-
llvm/test/CodeGen/NVPTX/local-stack-frame.ll | 67 +++--
llvm/test/CodeGen/NVPTX/lower-alloca.ll | 22 +-
.../CodeGen/NVPTX/lower-args-gridconstant.ll | 9 +-
llvm/test/CodeGen/NVPTX/lower-byval-args.ll | 231 ++++++++----------
llvm/test/CodeGen/NVPTX/vaargs.ll | 22 +-
llvm/test/CodeGen/NVPTX/variadics-backend.ll | 109 ++++-----
.../DebugInfo/NVPTX/dbg-declare-alloca.ll | 3 +-
.../DebugInfo/NVPTX/dbg-value-const-byref.ll | 2 +-
22 files changed, 342 insertions(+), 357 deletions(-)
diff --git a/llvm/include/llvm/Target/TargetMachine.h
b/llvm/include/llvm/Target/TargetMachine.h
index bf4e490554723..f4cd06bf93d40 100644
--- a/llvm/include/llvm/Target/TargetMachine.h
+++ b/llvm/include/llvm/Target/TargetMachine.h
@@ -208,7 +208,7 @@ class LLVM_ABI TargetMachine {
/// The LLVM Module owns a DataLayout that is used for the target independent
/// optimizations and code generation. This hook provides a target specific
/// check on the validity of this DataLayout.
- bool isCompatibleDataLayout(const DataLayout &Candidate) const {
+ virtual bool isCompatibleDataLayout(const DataLayout &Candidate) const {
return DL == Candidate;
}
diff --git a/llvm/lib/IR/Verifier.cpp b/llvm/lib/IR/Verifier.cpp
index 9d9b51db98702..2afa2bda53b52 100644
--- a/llvm/lib/IR/Verifier.cpp
+++ b/llvm/lib/IR/Verifier.cpp
@@ -114,6 +114,7 @@
#include "llvm/Pass.h"
#include "llvm/ProfileData/InstrProf.h"
#include "llvm/Support/AMDGPUAddrSpace.h"
+#include "llvm/Support/NVPTXAddrSpace.h"
#include "llvm/Support/AtomicOrdering.h"
#include "llvm/Support/Casting.h"
#include "llvm/Support/CommandLine.h"
@@ -4498,6 +4499,13 @@ void Verifier::visitAllocaInst(AllocaInst &AI) {
"alloca on amdgpu must be in addrspace(5)", &AI);
}
+ if (TT.isNVPTX()) {
+ Check(AI.getAddressSpace() == NVPTXAS::ADDRESS_SPACE_LOCAL ||
+ AI.getAddressSpace() == NVPTXAS::ADDRESS_SPACE_GENERIC,
+ "AllocaInst can only be in Generic or Local address space for
NVPTX.",
+ &AI);
+ }
+
visitInstruction(AI);
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index 7391c2d488b57..8528cb41d9244 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -1484,7 +1484,8 @@ void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters(
if (NumBytes) {
O << "\t.local .align " << MFI.getMaxAlign().value() << " .b8 \t"
<< DEPOTNAME << getFunctionNumber() << "[" << NumBytes << "];\n";
- if (static_cast<const NVPTXTargetMachine &>(MF.getTarget()).is64Bit()) {
+ if (static_cast<const NVPTXTargetMachine &>(MF.getTarget())
+ .getPointerSize(ADDRESS_SPACE_LOCAL) == 8) {
O << "\t.reg .b64 \t%SP;\n"
<< "\t.reg .b64 \t%SPL;\n";
} else {
diff --git a/llvm/lib/Target/NVPTX/NVPTXFrameLowering.cpp
b/llvm/lib/Target/NVPTX/NVPTXFrameLowering.cpp
index 47bc15f52bb96..78f18a93b869b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXFrameLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXFrameLowering.cpp
@@ -48,8 +48,8 @@ void NVPTXFrameLowering::emitPrologue(MachineFunction &MF,
// mov %SPL, %depot;
// cvta.local %SP, %SPL;
// for local address accesses in MF.
- bool Is64Bit =
- static_cast<const NVPTXTargetMachine &>(MF.getTarget()).is64Bit();
+ bool Is64Bit = static_cast<const NVPTXTargetMachine &>(MF.getTarget())
+ .getPointerSize(NVPTXAS::ADDRESS_SPACE_LOCAL) == 8;
unsigned CvtaLocalOpcode =
(Is64Bit ? NVPTX::cvta_local_64 : NVPTX::cvta_local);
unsigned MovDepotOpcode =
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index bb4bb1195f78b..f35407ed99106 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1803,10 +1803,7 @@ 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);
+ return Alloc;
}
SDValue NVPTXTargetLowering::LowerSTACKRESTORE(SDValue Op,
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
index 88bc000f39bf7..03412edb9e23c 100644
--- a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
@@ -6,16 +6,16 @@
//
//===----------------------------------------------------------------------===//
//
-// For all alloca instructions, and add a pair of cast to local address for
-// each of them. For example,
+// Change the Module's DataLayout to have the local address space for alloca's.
+// Change the address space of each alloca to local and add an addrspacecast to
+// generic address space. 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)*
+// %A = alloca i32, addrspace(5)
// %Generic = addrspacecast i32 addrspace(5)* %A to i32*
// store i32 0, i32 addrspace(5)* %Generic ; emits st.local.u32
//
@@ -24,18 +24,24 @@
//
//===----------------------------------------------------------------------===//
-#include "MCTargetDesc/NVPTXBaseInfo.h"
+#include "llvm/Support/NVPTXAddrSpace.h"
#include "NVPTX.h"
+#include "llvm/ADT/SmallVector.h"
#include "llvm/IR/Function.h"
+#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/Instructions.h"
+#include "llvm/IR/IntrinsicInst.h"
+#include "llvm/IR/Module.h"
#include "llvm/IR/Type.h"
#include "llvm/Pass.h"
using namespace llvm;
+using namespace NVPTXAS;
namespace {
class NVPTXLowerAlloca : public FunctionPass {
bool runOnFunction(Function &F) override;
+ bool doInitialization(Module &M) override;
public:
static char ID; // Pass identification, replacement for typeid
@@ -58,77 +64,55 @@ bool NVPTXLowerAlloca::runOnFunction(Function &F) {
if (skipFunction(F))
return false;
- bool Changed = false;
+ SmallVector<AllocaInst *, 16> Allocas;
for (auto &BB : F)
- for (auto &I : BB) {
- if (auto allocaInst = dyn_cast<AllocaInst>(&I)) {
- Changed = true;
+ for (auto &I : BB)
+ if (auto *Alloca = dyn_cast<AllocaInst>(&I))
+ if (Alloca->getAddressSpace() != ADDRESS_SPACE_LOCAL)
+ Allocas.push_back(Alloca);
+
+ if (Allocas.empty())
+ return false;
- PointerType *AllocInstPtrTy =
- cast<PointerType>(allocaInst->getType()->getScalarType());
- unsigned AllocAddrSpace = AllocInstPtrTy->getAddressSpace();
- assert((AllocAddrSpace == ADDRESS_SPACE_GENERIC ||
- AllocAddrSpace == ADDRESS_SPACE_LOCAL) &&
- "AllocaInst can only be in Generic or Local address space for "
- "NVPTX.");
+ for (AllocaInst *Alloca : Allocas) {
+ auto *NewAlloca = new AllocaInst(
+ Alloca->getAllocatedType(), ADDRESS_SPACE_LOCAL,
Alloca->getArraySize(),
+ Alloca->getAlign(), Alloca->getName());
+ auto *Cast = new AddrSpaceCastInst(
+ NewAlloca,
+ PointerType::get(Alloca->getAllocatedType()->getContext(),
+ ADDRESS_SPACE_GENERIC),
+ "");
+ Cast->insertBefore(Alloca->getIterator());
+ NewAlloca->insertBefore(Cast->getIterator());
+ for (auto &U : llvm::make_early_inc_range(Alloca->uses())) {
+ auto *II = dyn_cast<IntrinsicInst>(U.getUser());
+ if (!II || (II->getIntrinsicID() != Intrinsic::lifetime_start &&
+ II->getIntrinsicID() != Intrinsic::lifetime_end))
+ continue;
- Instruction *AllocaInLocalAS = allocaInst;
- auto ETy = allocaInst->getAllocatedType();
+ IRBuilder<> Builder(II);
+ Builder.CreateIntrinsic(II->getIntrinsicID(), {NewAlloca->getType()},
+ {NewAlloca});
+ II->eraseFromParent();
+ }
- // 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_LOCAL, 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(ETy->getContext(), ADDRESS_SPACE_LOCAL), "");
- ASCastToLocalAS->insertAfter(allocaInst->getIterator());
- AllocaInLocalAS = ASCastToLocalAS;
- }
+ Alloca->replaceAllUsesWith(Cast);
+ Alloca->eraseFromParent();
+ }
+ return true;
+}
- auto AllocaInGenericAS = new AddrSpaceCastInst(
- AllocaInLocalAS,
- PointerType::get(ETy->getContext(), ADDRESS_SPACE_GENERIC), "");
- AllocaInGenericAS->insertAfter(AllocaInLocalAS->getIterator());
+bool NVPTXLowerAlloca::doInitialization(Module &M) {
+ const auto &DL = M.getDataLayout();
+ if (DL.getAllocaAddrSpace() == ADDRESS_SPACE_LOCAL)
+ return false;
+ auto DLStr = DL.getStringRepresentation();
- 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;
+ auto AddrSpaceStr = "A" + std::to_string(ADDRESS_SPACE_LOCAL);
+ assert(!StringRef(DLStr).contains("A") && "DataLayout should not contain A");
+ M.setDataLayout(DLStr.empty() ? AddrSpaceStr : DLStr + "-" + AddrSpaceStr);
+ return true;
}
FunctionPass *llvm::createNVPTXLowerAllocaPass() {
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp
b/llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp
index e2bbe57c0085c..2b18ca9dd774a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp
@@ -529,7 +529,8 @@ void copyByValParam(Function &F, Argument &Arg) {
// the use of the byval parameter with this alloca instruction.
AllocA->setAlignment(
Arg.getParamAlign().value_or(DL.getPrefTypeAlign(StructType)));
- Arg.replaceAllUsesWith(AllocA);
+ auto *AddressSpaceCast = IRB.CreateAddrSpaceCast(AllocA, Arg.getType(),
Arg.getName());
+ Arg.replaceAllUsesWith(AddressSpaceCast);
CallInst *ArgInParam = createNVVMInternalAddrspaceWrap(IRB, Arg);
diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
index 646b554878c70..0c56caeadcffe 100644
--- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
@@ -119,7 +119,7 @@ bool
NVPTXRegisterInfo::eliminateFrameIndex(MachineBasicBlock::iterator II,
MI.getOperand(FIOperandNum + 1).getImm();
// Using I0 as the frame pointer
- MI.getOperand(FIOperandNum).ChangeToRegister(getFrameRegister(MF), false);
+ MI.getOperand(FIOperandNum).ChangeToRegister(getFrameLocalRegister(MF),
false);
MI.getOperand(FIOperandNum + 1).ChangeToImmediate(Offset);
return false;
}
@@ -127,14 +127,18 @@ bool
NVPTXRegisterInfo::eliminateFrameIndex(MachineBasicBlock::iterator II,
Register NVPTXRegisterInfo::getFrameRegister(const MachineFunction &MF) const {
const NVPTXTargetMachine &TM =
static_cast<const NVPTXTargetMachine &>(MF.getTarget());
- return TM.is64Bit() ? NVPTX::VRFrame64 : NVPTX::VRFrame32;
+ return TM.getPointerSize(NVPTXAS::ADDRESS_SPACE_LOCAL) == 8
+ ? NVPTX::VRFrame64
+ : NVPTX::VRFrame32;
}
Register
NVPTXRegisterInfo::getFrameLocalRegister(const MachineFunction &MF) const {
const NVPTXTargetMachine &TM =
static_cast<const NVPTXTargetMachine &>(MF.getTarget());
- return TM.is64Bit() ? NVPTX::VRFrameLocal64 : NVPTX::VRFrameLocal32;
+ return TM.getPointerSize(NVPTXAS::ADDRESS_SPACE_LOCAL) == 8
+ ? NVPTX::VRFrameLocal64
+ : NVPTX::VRFrameLocal32;
}
void NVPTXRegisterInfo::clearDebugRegisterMap() const {
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
index 0603994606d71..209dd51a44a61 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
@@ -371,6 +371,8 @@ void NVPTXPassConfig::addIRPasses() {
if (getOptLevel() != CodeGenOptLevel::None) {
addAddressSpaceInferencePasses();
addStraightLineScalarOptimizationPasses();
+ } else {
+ addPass(createNVPTXLowerAllocaPass());
}
addPass(createAtomicExpandLegacyPass());
@@ -502,3 +504,25 @@ void NVPTXPassConfig::addMachineSSAOptimization() {
addPass(&PeepholeOptimizerLegacyID);
printAndVerify("After codegen peephole optimization pass");
}
+
+bool NVPTXTargetMachine::isCompatibleDataLayout(
+ const DataLayout &Candidate) const {
+ //XXX: Should we enforce that the Candidate DataLayout has the same address
space for allocas?
+ if (DL == Candidate)
+ return true;
+
+ auto DLStr = DL.getStringRepresentation();
+ if (!StringRef(DLStr).contains("A"))
+ DLStr = DLStr.empty() ? "A" + std::to_string(ADDRESS_SPACE_LOCAL)
+ : DLStr + "-A" + std::to_string(ADDRESS_SPACE_LOCAL);
+ auto NewDL = DataLayout(DLStr);
+
+ return NewDL == Candidate;
+}
+
+unsigned NVPTXTargetMachine::getAddressSpaceForPseudoSourceKind(unsigned Kind)
const {
+ if (Kind == PseudoSourceValue::FixedStack) {
+ return ADDRESS_SPACE_LOCAL;
+ }
+ return CodeGenTargetMachineImpl::getAddressSpaceForPseudoSourceKind(Kind);
+}
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.h
b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.h
index 118a01a0352f5..c2f09a89865c4 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.h
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.h
@@ -76,6 +76,9 @@ class NVPTXTargetMachine : public CodeGenTargetMachineImpl {
std::pair<const Value *, unsigned>
getPredicatedAddrSpace(const Value *V) const override;
+
+ bool isCompatibleDataLayout(const DataLayout &Candidate) const override;
+ unsigned getAddressSpaceForPseudoSourceKind(unsigned Kind) const override;
}; // NVPTXTargetMachine.
class NVPTXTargetMachine32 : public NVPTXTargetMachine {
diff --git a/llvm/test/CodeGen/NVPTX/call-with-alloca-buffer.ll
b/llvm/test/CodeGen/NVPTX/call-with-alloca-buffer.ll
index 0eb7f6462f6fa..e825f8ef19949 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[[SP_REG0:[0-9]+]], %SPL, 0
+; CHECK: cvta.local.u64 %rd[[SP_REG:[0-9]+]], %rd[[SP_REG0]];
; 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/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/f32x2-instructions.ll
b/llvm/test/CodeGen/NVPTX/f32x2-instructions.ll
index 7ca16f702d8f3..cf00c2c9eaa3f 100644
--- a/llvm/test/CodeGen/NVPTX/f32x2-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/f32x2-instructions.ll
@@ -86,15 +86,14 @@ define float @test_extract_i(<2 x float> %a, i64 %idx) #0 {
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
; CHECK-NOF32X2-NEXT: mov.b64 %SPL, __local_depot3;
-; CHECK-NOF32X2-NEXT: cvta.local.u64 %SP, %SPL;
; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_extract_i_param_0];
; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_extract_i_param_1];
-; CHECK-NOF32X2-NEXT: st.v2.b32 [%SP], {%r1, %r2};
+; CHECK-NOF32X2-NEXT: st.local.v2.b32 [%SPL], {%r1, %r2};
; CHECK-NOF32X2-NEXT: and.b64 %rd2, %rd1, 1;
; CHECK-NOF32X2-NEXT: shl.b64 %rd3, %rd2, 2;
-; CHECK-NOF32X2-NEXT: add.u64 %rd4, %SP, 0;
+; CHECK-NOF32X2-NEXT: add.u64 %rd4, %SPL, 0;
; CHECK-NOF32X2-NEXT: or.b64 %rd5, %rd4, %rd3;
-; CHECK-NOF32X2-NEXT: ld.b32 %r3, [%rd5];
+; CHECK-NOF32X2-NEXT: ld.local.b32 %r3, [%rd5];
; CHECK-NOF32X2-NEXT: st.param.b32 [func_retval0], %r3;
; CHECK-NOF32X2-NEXT: ret;
;
diff --git a/llvm/test/CodeGen/NVPTX/indirect_byval.ll
b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
index 673fb73948268..39813efef9b64 100644
--- a/llvm/test/CodeGen/NVPTX/indirect_byval.ll
+++ b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
@@ -21,19 +21,18 @@ define internal i32 @foo() {
; CHECK-EMPTY:
; 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.local.b8 %rs1, [%SPL+1];
; CHECK-NEXT: st.param.b8 [param0], %rs1;
; CHECK-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .align
1 .b8 _[1], .param .b64 _);
-; 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;
@@ -58,19 +57,18 @@ define internal i32 @bar() {
; CHECK-EMPTY:
; 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.local.b64 %rd4, [%SPL+8];
; CHECK-NEXT: st.param.b64 [param0], %rd4;
; CHECK-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .align
8 .b8 _[8], .param .b64 _);
-; 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 f7137e05a5e4f..3e6f5e37aa85f 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;
;
@@ -139,21 +134,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;
@@ -167,21 +161,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 57c1e5826c89a..7bc90c42ad12f 100644
--- a/llvm/test/CodeGen/NVPTX/lower-alloca.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-alloca.ll
@@ -7,28 +7,28 @@ target datalayout =
"e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f3
target triple = "nvptx64-unknown-unknown"
define ptx_kernel void @kernel() {
-; LABEL: @lower_alloca
+; CHECK-LABEL: @kernel
; 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
+; CHECK: %A1 = alloca i32, align 4, addrspace(5)
+; CHECK: store i32 0, ptr addrspace(5) %A1
+; LOWERALLOCAONLY: %A1 = alloca i32, align 4, addrspace(5)
+; LOWERALLOCAONLY: [[V2:%.*]] = addrspacecast ptr addrspace(5) %A1 to ptr
; LOWERALLOCAONLY: store i32 0, ptr [[V2]], align 4
-; PTX: st.local.b32 [{{%rd[0-9]+}}], 0
+; PTX: st.local.b32 [%SPL], 0
store i32 0, ptr %A
call void @callee(ptr %A)
ret void
}
define void @alloca_in_explicit_local_as() {
-; LABEL: @lower_alloca_addrspace5
+; CHECK-LABEL: @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
-; LOWERALLOCAONLY: [[V1:%.*]] = addrspacecast ptr addrspace(5) %A to ptr
-; LOWERALLOCAONLY: store i32 0, ptr [[V1]], align 4
+; CHECK: store i32 0, ptr addrspace(5) %A, align 4
+; PTX: st.local.b32 [%SPL], 0
+; LOWERALLOCAONLY: %A = alloca i32, align 4, addrspace(5)
+; LOWERALLOCAONLY: store i32 0, ptr addrspace(5) %A
store i32 0, ptr addrspace(5) %A
call void @callee(ptr addrspace(5) %A)
ret void
diff --git a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
index 8adde4ceefbf4..f4a24ab41ba17 100644
--- a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
@@ -156,22 +156,21 @@ define ptx_kernel void @multiple_grid_const_escape(ptr
byval(%struct.s) align 4
; PTX-EMPTY:
; 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: ld.param.b32 %r1, [multiple_grid_const_escape_param_1];
; PTX-NEXT: mov.b64 %rd2, multiple_grid_const_escape_param_2;
; PTX-NEXT: cvta.param.u64 %rd3, %rd2;
; 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: prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _,
.param .b64 _, .param .b64 _);
; PTX-NEXT: mov.b64 %rd7, escape3;
diff --git a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
index 21257e21bea9f..d8a9530e931a6 100644
--- a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
@@ -135,24 +135,23 @@ 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: 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 [%SPL], %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;
@@ -175,25 +174,24 @@ 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: 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], %rd3;
+; PTX-NEXT: st.param.b64 [param0], %rd7;
; PTX-NEXT: call.uni _Z6escapePv, (param0);
; PTX-NEXT: } // callseq 1
; PTX-NEXT: ret;
@@ -216,24 +214,23 @@ 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;
-; 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 %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: 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 [%SPL], %rd8;
+; PTX-NEXT: st.global.b64 [%rd2], %rd4;
; PTX-NEXT: ret;
entry:
store ptr %s, ptr %out, align 8
@@ -254,25 +251,24 @@ 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: 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:
%b = getelementptr inbounds nuw i8, ptr %s, i64 4
@@ -294,24 +290,23 @@ 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;
-; 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 %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: 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 [%SPL], %rd8;
+; PTX-NEXT: st.global.b64 [%rd2], %rd4;
; PTX-NEXT: ret;
entry:
%i = ptrtoint ptr %s to i64
@@ -455,64 +450,51 @@ define dso_local ptx_kernel void @memcpy_to_param(ptr
nocapture noundef readonly
; PTX-NEXT: .local .align 8 .b8 __local_depot9[8];
; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
-; PTX-NEXT: .reg .b32 %r<3>;
-; PTX-NEXT: .reg .b64 %rd<47>;
+; PTX-NEXT: .reg .b16 %rs<17>;
+; PTX-NEXT: .reg .b64 %rd<8>;
; 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: add.u64 %rd2, %SPL, 0;
-; PTX-NEXT: ld.param.b32 %r1, [memcpy_to_param_param_1+4];
-; PTX-NEXT: st.local.b32 [%rd2+4], %r1;
-; PTX-NEXT: ld.param.b32 %r2, [memcpy_to_param_param_1];
-; PTX-NEXT: st.local.b32 [%rd2], %r2;
-; PTX-NEXT: ld.volatile.b8 %rd3, [%rd1];
-; PTX-NEXT: ld.volatile.b8 %rd4, [%rd1+1];
-; PTX-NEXT: shl.b64 %rd5, %rd4, 8;
-; PTX-NEXT: or.b64 %rd6, %rd5, %rd3;
-; PTX-NEXT: ld.volatile.b8 %rd7, [%rd1+2];
-; PTX-NEXT: shl.b64 %rd8, %rd7, 16;
-; PTX-NEXT: ld.volatile.b8 %rd9, [%rd1+3];
-; PTX-NEXT: shl.b64 %rd10, %rd9, 24;
-; PTX-NEXT: or.b64 %rd11, %rd10, %rd8;
-; PTX-NEXT: or.b64 %rd12, %rd11, %rd6;
-; PTX-NEXT: ld.volatile.b8 %rd13, [%rd1+4];
-; PTX-NEXT: ld.volatile.b8 %rd14, [%rd1+5];
-; PTX-NEXT: shl.b64 %rd15, %rd14, 8;
-; PTX-NEXT: or.b64 %rd16, %rd15, %rd13;
-; PTX-NEXT: ld.volatile.b8 %rd17, [%rd1+6];
-; PTX-NEXT: shl.b64 %rd18, %rd17, 16;
-; PTX-NEXT: ld.volatile.b8 %rd19, [%rd1+7];
-; PTX-NEXT: shl.b64 %rd20, %rd19, 24;
-; PTX-NEXT: or.b64 %rd21, %rd20, %rd18;
-; PTX-NEXT: or.b64 %rd22, %rd21, %rd16;
-; PTX-NEXT: shl.b64 %rd23, %rd22, 32;
-; PTX-NEXT: or.b64 %rd24, %rd23, %rd12;
-; PTX-NEXT: st.volatile.b64 [%SP], %rd24;
-; PTX-NEXT: ld.volatile.b8 %rd25, [%rd1+8];
-; PTX-NEXT: ld.volatile.b8 %rd26, [%rd1+9];
-; PTX-NEXT: shl.b64 %rd27, %rd26, 8;
-; PTX-NEXT: or.b64 %rd28, %rd27, %rd25;
-; PTX-NEXT: ld.volatile.b8 %rd29, [%rd1+10];
-; PTX-NEXT: shl.b64 %rd30, %rd29, 16;
-; PTX-NEXT: ld.volatile.b8 %rd31, [%rd1+11];
-; PTX-NEXT: shl.b64 %rd32, %rd31, 24;
-; PTX-NEXT: or.b64 %rd33, %rd32, %rd30;
-; PTX-NEXT: or.b64 %rd34, %rd33, %rd28;
-; PTX-NEXT: ld.volatile.b8 %rd35, [%rd1+12];
-; PTX-NEXT: ld.volatile.b8 %rd36, [%rd1+13];
-; PTX-NEXT: shl.b64 %rd37, %rd36, 8;
-; PTX-NEXT: or.b64 %rd38, %rd37, %rd35;
-; PTX-NEXT: ld.volatile.b8 %rd39, [%rd1+14];
-; PTX-NEXT: shl.b64 %rd40, %rd39, 16;
-; PTX-NEXT: ld.volatile.b8 %rd41, [%rd1+15];
-; PTX-NEXT: shl.b64 %rd42, %rd41, 24;
-; PTX-NEXT: or.b64 %rd43, %rd42, %rd40;
-; PTX-NEXT: or.b64 %rd44, %rd43, %rd38;
-; PTX-NEXT: shl.b64 %rd45, %rd44, 32;
-; PTX-NEXT: or.b64 %rd46, %rd45, %rd34;
-; PTX-NEXT: st.volatile.b64 [%SP+8], %rd46;
+; PTX-NEXT: cvta.local.u64 %rd3, %rd2;
+; PTX-NEXT: ld.param.b32 %rd4, [memcpy_to_param_param_1+4];
+; PTX-NEXT: shl.b64 %rd5, %rd4, 32;
+; PTX-NEXT: ld.param.b32 %rd6, [memcpy_to_param_param_1];
+; PTX-NEXT: or.b64 %rd7, %rd5, %rd6;
+; PTX-NEXT: st.local.b64 [%SPL], %rd7;
+; PTX-NEXT: ld.volatile.b8 %rs1, [%rd1];
+; PTX-NEXT: st.volatile.b8 [%rd3], %rs1;
+; PTX-NEXT: ld.volatile.b8 %rs2, [%rd1+1];
+; PTX-NEXT: st.volatile.b8 [%rd3+1], %rs2;
+; PTX-NEXT: ld.volatile.b8 %rs3, [%rd1+2];
+; PTX-NEXT: st.volatile.b8 [%rd3+2], %rs3;
+; PTX-NEXT: ld.volatile.b8 %rs4, [%rd1+3];
+; PTX-NEXT: st.volatile.b8 [%rd3+3], %rs4;
+; PTX-NEXT: ld.volatile.b8 %rs5, [%rd1+4];
+; PTX-NEXT: st.volatile.b8 [%rd3+4], %rs5;
+; PTX-NEXT: ld.volatile.b8 %rs6, [%rd1+5];
+; PTX-NEXT: st.volatile.b8 [%rd3+5], %rs6;
+; PTX-NEXT: ld.volatile.b8 %rs7, [%rd1+6];
+; PTX-NEXT: st.volatile.b8 [%rd3+6], %rs7;
+; PTX-NEXT: ld.volatile.b8 %rs8, [%rd1+7];
+; PTX-NEXT: st.volatile.b8 [%rd3+7], %rs8;
+; PTX-NEXT: ld.volatile.b8 %rs9, [%rd1+8];
+; PTX-NEXT: st.volatile.b8 [%rd3+8], %rs9;
+; PTX-NEXT: ld.volatile.b8 %rs10, [%rd1+9];
+; PTX-NEXT: st.volatile.b8 [%rd3+9], %rs10;
+; PTX-NEXT: ld.volatile.b8 %rs11, [%rd1+10];
+; PTX-NEXT: st.volatile.b8 [%rd3+10], %rs11;
+; PTX-NEXT: ld.volatile.b8 %rs12, [%rd1+11];
+; PTX-NEXT: st.volatile.b8 [%rd3+11], %rs12;
+; PTX-NEXT: ld.volatile.b8 %rs13, [%rd1+12];
+; PTX-NEXT: st.volatile.b8 [%rd3+12], %rs13;
+; PTX-NEXT: ld.volatile.b8 %rs14, [%rd1+13];
+; PTX-NEXT: st.volatile.b8 [%rd3+13], %rs14;
+; PTX-NEXT: ld.volatile.b8 %rs15, [%rd1+14];
+; PTX-NEXT: st.volatile.b8 [%rd3+14], %rs15;
+; PTX-NEXT: ld.volatile.b8 %rs16, [%rd1+15];
+; PTX-NEXT: st.volatile.b8 [%rd3+15], %rs16;
; PTX-NEXT: ret;
entry:
tail call void @llvm.memcpy.p0.p0.i64(ptr %s, ptr %in, i64 16, i1 true)
@@ -655,17 +637,16 @@ 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.b32 [%SP], %r1;
+; PTX-NEXT: st.local.b32 [%SPL], %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 [%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;
bb:
@@ -830,25 +811,23 @@ 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;
-; 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: 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 [%SPL], %r1;
; PTX-NEXT: ld.param.b32 %r2, [test_phi_write_param_0];
-; PTX-NEXT: st.b32 [%SP+4], %r2;
+; 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: mov.b64 %rd2, %rd1;
+; PTX-NEXT: add.u64 %rd1, %SPL, 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
@@ -882,13 +861,11 @@ define ptx_kernel void @test_forward_byval_arg(ptr
byval(i32) align 4 %input) {
; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
; PTX-NEXT: .reg .b32 %r<2>;
-; PTX-NEXT: .reg .b64 %rd<2>;
; PTX-EMPTY:
; PTX-NEXT: // %bb.0:
; PTX-NEXT: mov.b64 %SPL, __local_depot15;
-; PTX-NEXT: add.u64 %rd1, %SPL, 0;
; PTX-NEXT: ld.param.b32 %r1, [test_forward_byval_arg_param_0];
-; PTX-NEXT: st.local.b32 [%rd1], %r1;
+; PTX-NEXT: st.local.b32 [%SPL], %r1;
; PTX-NEXT: { // callseq 2, 0
; PTX-NEXT: .param .align 4 .b8 param0[4];
; PTX-NEXT: st.param.b32 [param0], %r1;
diff --git a/llvm/test/CodeGen/NVPTX/vaargs.ll
b/llvm/test/CodeGen/NVPTX/vaargs.ll
index 9e312a2fec60a..19ed05e28a3d4 100644
--- a/llvm/test/CodeGen/NVPTX/vaargs.ll
+++ b/llvm/test/CodeGen/NVPTX/vaargs.ll
@@ -17,55 +17,55 @@ entry:
; Test va_start
; CHECK: .param .align 8 .b8 foo_vararg[]
; CHECK: mov.b[[BITS]] [[VA_PTR:%(r|rd)[0-9]+]], foo_vararg;
-; CHECK-NEXT: st.b[[BITS]] [%SP], [[VA_PTR]];
+; CHECK-NEXT: st.b[[BITS]] [[[SP1:%(r|rd)[0-9]+]]], [[VA_PTR]]
call void @llvm.va_start(ptr %al)
; Test va_copy()
-; CHECK-NEXT: ld.b[[BITS]] [[VA_PTR:%(r|rd)[0-9]+]], [%SP];
-; CHECK-NEXT: st.b[[BITS]] [%SP+{{[0-9]+}}], [[VA_PTR]];
+; CHECK-NEXT: ld.b[[BITS]] [[VA_PTR:%(r|rd)[0-9]+]], [[[SP1]]];
+; CHECK-NEXT: st.b[[BITS]] [[[SP2:%(r|rd)[0-9]+]]], [[VA_PTR]];
call void @llvm.va_copy(ptr %al2, ptr %al)
; Test va_arg(ap, int32_t)
-; CHECK-NEXT: ld.b[[BITS]] [[VA_PTR:%(r|rd)[0-9]+]], [%SP];
+; CHECK-NEXT: ld.b[[BITS]] [[VA_PTR:%(r|rd)[0-9]+]], [[[SP1]]];
; CHECK-NEXT: add.s[[BITS]] [[VA_PTR_TMP:%(r|rd)[0-9]+]], [[VA_PTR]], 3;
; CHECK-NEXT: and.b[[BITS]] [[VA_PTR_ALIGN:%(r|rd)[0-9]+]], [[VA_PTR_TMP]],
-4;
; CHECK-NEXT: add.s[[BITS]] [[VA_PTR_NEXT:%(r|rd)[0-9]+]],
[[VA_PTR_ALIGN]], 4;
-; CHECK-NEXT: st.b[[BITS]] [%SP], [[VA_PTR_NEXT]];
+; CHECK-NEXT: st.b[[BITS]] [[[SP1]]], [[VA_PTR_NEXT]];
; CHECK-NEXT: ld.local.b32 %r{{[0-9]+}}, [[[VA_PTR_ALIGN]]];
%0 = va_arg ptr %al, i32
; Test va_arg(ap, int64_t)
-; CHECK-NEXT: ld.b[[BITS]] [[VA_PTR:%(r|rd)[0-9]+]], [%SP];
+; CHECK-NEXT: ld.b[[BITS]] [[VA_PTR:%(r|rd)[0-9]+]], [[[SP1]]];
; CHECK-NEXT: add.s[[BITS]] [[VA_PTR_TMP:%(r|rd)[0-9]+]], [[VA_PTR]], 7;
; CHECK-NEXT: and.b[[BITS]] [[VA_PTR_ALIGN:%(r|rd)[0-9]+]], [[VA_PTR_TMP]],
-8;
; CHECK-NEXT: add.s[[BITS]] [[VA_PTR_NEXT:%(r|rd)[0-9]+]],
[[VA_PTR_ALIGN]], 8;
-; CHECK-NEXT: st.b[[BITS]] [%SP], [[VA_PTR_NEXT]];
+; CHECK-NEXT: st.b[[BITS]] [[[SP1]]], [[VA_PTR_NEXT]];
; CHECK-NEXT: ld.local.b64 %rd{{[0-9]+}}, [[[VA_PTR_ALIGN]]];
%1 = va_arg ptr %al, i64
; Test va_arg(ap, double)
-; CHECK-NEXT: ld.b[[BITS]] [[VA_PTR:%(r|rd)[0-9]+]], [%SP];
+; CHECK-NEXT: ld.b[[BITS]] [[VA_PTR:%(r|rd)[0-9]+]], [[[SP1]]];
; CHECK-NEXT: add.s[[BITS]] [[VA_PTR_TMP:%(r|rd)[0-9]+]], [[VA_PTR]], 7;
; CHECK-NEXT: and.b[[BITS]] [[VA_PTR_ALIGN:%(r|rd)[0-9]+]], [[VA_PTR_TMP]],
-8;
; CHECK-NEXT: add.s[[BITS]] [[VA_PTR_NEXT:%(r|rd)[0-9]+]],
[[VA_PTR_ALIGN]], 8;
-; CHECK-NEXT: st.b[[BITS]] [%SP], [[VA_PTR_NEXT]];
+; CHECK-NEXT: st.b[[BITS]] [[[SP1]]], [[VA_PTR_NEXT]];
; CHECK-NEXT: ld.local.b64 %rd{{[0-9]+}}, [[[VA_PTR_ALIGN]]];
%2 = va_arg ptr %al, double
; Test va_arg(ap, ptr)
-; CHECK-NEXT: ld.b[[BITS]] [[VA_PTR:%(r|rd)[0-9]+]], [%SP];
+; CHECK-NEXT: ld.b[[BITS]] [[VA_PTR:%(r|rd)[0-9]+]], [[[SP1]]];
; CHECK32-NEXT: add.s32 [[VA_PTR_TMP:%r[0-9]+]], [[VA_PTR]], 3;
; CHECK64-NEXT: add.s64 [[VA_PTR_TMP:%rd[0-9]+]], [[VA_PTR]], 7;
; CHECK32-NEXT: and.b32 [[VA_PTR_ALIGN:%r[0-9]+]], [[VA_PTR_TMP]], -4;
; CHECK64-NEXT: and.b64 [[VA_PTR_ALIGN:%rd[0-9]+]], [[VA_PTR_TMP]], -8;
; CHECK32-NEXT: add.s32 [[VA_PTR_NEXT:%r[0-9]+]], [[VA_PTR_ALIGN]], 4;
; CHECK64-NEXT: add.s64 [[VA_PTR_NEXT:%rd[0-9]+]], [[VA_PTR_ALIGN]], 8;
-; CHECK-NEXT: st.b[[BITS]] [%SP], [[VA_PTR_NEXT]];
+; CHECK-NEXT: st.b[[BITS]] [[[SP1]]], [[VA_PTR_NEXT]];
; CHECK-NEXT: ld.local.b[[BITS]] %{{(r|rd)[0-9]+}}, [[[VA_PTR_ALIGN]]];
%3 = va_arg ptr %al, ptr
diff --git a/llvm/test/CodeGen/NVPTX/variadics-backend.ll
b/llvm/test/CodeGen/NVPTX/variadics-backend.ll
index 890753b6ac5aa..7fe0dda2bf94e 100644
--- a/llvm/test/CodeGen/NVPTX/variadics-backend.ll
+++ b/llvm/test/CodeGen/NVPTX/variadics-backend.ll
@@ -105,22 +105,22 @@ define dso_local i32 @foo() {
; CHECK-PTX-NEXT: .reg .b64 %SP;
; CHECK-PTX-NEXT: .reg .b64 %SPL;
; CHECK-PTX-NEXT: .reg .b32 %r<2>;
-; CHECK-PTX-NEXT: .reg .b64 %rd<2>;
+; CHECK-PTX-NEXT: .reg .b64 %rd<3>;
; 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.b64 [%SP], 4294967297;
-; CHECK-PTX-NEXT: st.b32 [%SP+8], 1;
-; CHECK-PTX-NEXT: st.b64 [%SP+16], 1;
-; CHECK-PTX-NEXT: st.b64 [%SP+24], 4607182418800017408;
-; CHECK-PTX-NEXT: st.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: add.u64 %rd1, %SPL, 0;
+; CHECK-PTX-NEXT: cvta.local.u64 %rd2, %rd1;
; 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.b64 [param1], %rd2;
; CHECK-PTX-NEXT: st.param.b32 [param0], 1;
; CHECK-PTX-NEXT: call.uni (retval0), variadics1, (param0, param1);
; CHECK-PTX-NEXT: ld.param.b32 %r1, [retval0];
@@ -138,34 +138,34 @@ 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: 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.b32 %r2, [%rd4];
-; CHECK-PTX-NEXT: ld.s8 %r3, [%rd4+4];
-; CHECK-PTX-NEXT: ld.b8 %rs1, [%rd4+7];
-; CHECK-PTX-NEXT: st.local.b8 [%rd2+2], %rs1;
-; CHECK-PTX-NEXT: ld.b8 %rs2, [%rd4+6];
-; CHECK-PTX-NEXT: st.local.b8 [%rd2+1], %rs2;
-; CHECK-PTX-NEXT: ld.b8 %rs3, [%rd4+5];
-; CHECK-PTX-NEXT: st.local.b8 [%rd2], %rs3;
-; CHECK-PTX-NEXT: ld.b64 %rd5, [%rd4+8];
+; CHECK-PTX-NEXT: add.s64 %rd2, %rd1, 7;
+; CHECK-PTX-NEXT: and.b64 %rd3, %rd2, -8;
+; CHECK-PTX-NEXT: ld.b32 %r2, [%rd3];
+; CHECK-PTX-NEXT: ld.s8 %r3, [%rd3+4];
+; CHECK-PTX-NEXT: ld.b8 %rs1, [%rd3+7];
+; CHECK-PTX-NEXT: st.local.b8 [%SPL+2], %rs1;
+; CHECK-PTX-NEXT: ld.b8 %rs2, [%rd3+5];
+; CHECK-PTX-NEXT: ld.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 [%SPL], %rs5;
+; CHECK-PTX-NEXT: ld.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 +201,28 @@ 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-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: add.u64 %rd1, %SPL, 0;
; 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 [%SPL+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.b32 [%SP+8], 1;
-; CHECK-PTX-NEXT: st.b8 [%SP+12], 1;
-; CHECK-PTX-NEXT: st.b64 [%SP+16], 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 [%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: add.u64 %rd1, %SPL, 8;
+; CHECK-PTX-NEXT: cvta.local.u64 %rd2, %rd1;
; 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 %rd2, %SP, 8;
; CHECK-PTX-NEXT: st.param.b64 [param1], %rd2;
; CHECK-PTX-NEXT: st.param.b32 [param0], 1;
; CHECK-PTX-NEXT: call.uni (retval0), variadics2, (param0, param1);
@@ -283,18 +283,18 @@ define dso_local i32 @baz() {
; CHECK-PTX-NEXT: .reg .b64 %SP;
; CHECK-PTX-NEXT: .reg .b64 %SPL;
; CHECK-PTX-NEXT: .reg .b32 %r<2>;
-; CHECK-PTX-NEXT: .reg .b64 %rd<2>;
+; CHECK-PTX-NEXT: .reg .b64 %rd<3>;
; 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.v4.b32 [%SP], {1, 1, 1, 1};
+; CHECK-PTX-NEXT: st.local.v4.b32 [%SPL], {1, 1, 1, 1};
+; CHECK-PTX-NEXT: add.u64 %rd1, %SPL, 0;
+; CHECK-PTX-NEXT: cvta.local.u64 %rd2, %rd1;
; 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.b64 [param1], %rd2;
; CHECK-PTX-NEXT: st.param.b32 [param0], 1;
; CHECK-PTX-NEXT: call.uni (retval0), variadics3, (param0, param1);
; CHECK-PTX-NEXT: ld.param.b32 %r1, [retval0];
@@ -348,27 +348,24 @@ 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<5>;
; 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: add.u64 %rd1, %SPL, 0;
-; CHECK-PTX-NEXT: ld.global.nc.b64 %rd2, [__const_$_qux_$_s+8];
-; CHECK-PTX-NEXT: st.local.b64 [%rd1+8], %rd2;
-; CHECK-PTX-NEXT: ld.global.nc.b64 %rd3, [__const_$_qux_$_s];
-; CHECK-PTX-NEXT: st.local.b64 [%rd1], %rd3;
-; CHECK-PTX-NEXT: st.b64 [%SP+16], 1;
+; CHECK-PTX-NEXT: ld.global.nc.b64 %rd1, [__const_$_qux_$_s+8];
+; 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 [%SPL], %rd2;
+; CHECK-PTX-NEXT: st.local.b64 [%SPL+16], 1;
+; CHECK-PTX-NEXT: add.u64 %rd3, %SPL, 16;
+; CHECK-PTX-NEXT: cvta.local.u64 %rd4, %rd3;
; 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 %rd4, %SP, 16;
; CHECK-PTX-NEXT: st.param.b64 [param1], %rd4;
-; 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: st.param.b64 [param0+8], %rd1;
+; CHECK-PTX-NEXT: st.param.b64 [param0], %rd2;
; CHECK-PTX-NEXT: call.uni (retval0), variadics4, (param0, param1);
; CHECK-PTX-NEXT: } // callseq 3
; CHECK-PTX-NEXT: ret;
diff --git a/llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll
b/llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll
index fa42481016540..5d9b4ee6f5297 100644
--- a/llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll
+++ b/llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll
@@ -9,8 +9,7 @@
; 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
diff --git a/llvm/test/DebugInfo/NVPTX/dbg-value-const-byref.ll
b/llvm/test/DebugInfo/NVPTX/dbg-value-const-byref.ll
index 41734f33213e8..ffab7a1b6cfc6 100644
--- a/llvm/test/DebugInfo/NVPTX/dbg-value-const-byref.ll
+++ b/llvm/test/DebugInfo/NVPTX/dbg-value-const-byref.ll
@@ -22,7 +22,7 @@
; CHECK: DEBUG_VALUE: foo:i <- 3
; CHECK: DEBUG_VALUE: foo:i <- 7
; CHECK: DEBUG_VALUE: foo:i <- %
-; CHECK: DEBUG_VALUE: foo:i <- [DW_OP_deref] $vrdepot
+; CHECK: DEBUG_VALUE: foo:i <- [DW_OP_deref] %
; Function Attrs: nounwind ssp uwtable
define i32 @foo() #0 !dbg !4 {
>From 7eb8b83a4a30550ee5ded74f29660f83d5dabc92 Mon Sep 17 00:00:00 2001
From: Theodoros Theodoridis <ttheodori...@nvidia.com>
Date: Thu, 21 Aug 2025 18:43:11 +0000
Subject: [PATCH 2/7] format
---
llvm/lib/IR/Verifier.cpp | 2 +-
llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp | 2 +-
llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp | 3 ++-
llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp | 3 ++-
llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp | 8 +++++---
5 files changed, 11 insertions(+), 7 deletions(-)
diff --git a/llvm/lib/IR/Verifier.cpp b/llvm/lib/IR/Verifier.cpp
index 2afa2bda53b52..2a4d4bfa1a4c7 100644
--- a/llvm/lib/IR/Verifier.cpp
+++ b/llvm/lib/IR/Verifier.cpp
@@ -114,13 +114,13 @@
#include "llvm/Pass.h"
#include "llvm/ProfileData/InstrProf.h"
#include "llvm/Support/AMDGPUAddrSpace.h"
-#include "llvm/Support/NVPTXAddrSpace.h"
#include "llvm/Support/AtomicOrdering.h"
#include "llvm/Support/Casting.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/MathExtras.h"
#include "llvm/Support/ModRef.h"
+#include "llvm/Support/NVPTXAddrSpace.h"
#include "llvm/Support/raw_ostream.h"
#include <algorithm>
#include <cassert>
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
index 03412edb9e23c..4477be9b65004 100644
--- a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
@@ -24,7 +24,6 @@
//
//===----------------------------------------------------------------------===//
-#include "llvm/Support/NVPTXAddrSpace.h"
#include "NVPTX.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/IR/Function.h"
@@ -34,6 +33,7 @@
#include "llvm/IR/Module.h"
#include "llvm/IR/Type.h"
#include "llvm/Pass.h"
+#include "llvm/Support/NVPTXAddrSpace.h"
using namespace llvm;
using namespace NVPTXAS;
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp
b/llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp
index 2b18ca9dd774a..050ba38b7fcd2 100644
--- a/llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXLowerArgs.cpp
@@ -529,7 +529,8 @@ void copyByValParam(Function &F, Argument &Arg) {
// the use of the byval parameter with this alloca instruction.
AllocA->setAlignment(
Arg.getParamAlign().value_or(DL.getPrefTypeAlign(StructType)));
- auto *AddressSpaceCast = IRB.CreateAddrSpaceCast(AllocA, Arg.getType(),
Arg.getName());
+ auto *AddressSpaceCast =
+ IRB.CreateAddrSpaceCast(AllocA, Arg.getType(), Arg.getName());
Arg.replaceAllUsesWith(AddressSpaceCast);
CallInst *ArgInParam = createNVVMInternalAddrspaceWrap(IRB, Arg);
diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
index 0c56caeadcffe..04eb095f31180 100644
--- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
@@ -119,7 +119,8 @@ bool
NVPTXRegisterInfo::eliminateFrameIndex(MachineBasicBlock::iterator II,
MI.getOperand(FIOperandNum + 1).getImm();
// Using I0 as the frame pointer
- MI.getOperand(FIOperandNum).ChangeToRegister(getFrameLocalRegister(MF),
false);
+ MI.getOperand(FIOperandNum)
+ .ChangeToRegister(getFrameLocalRegister(MF), 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 209dd51a44a61..0874aec69fb40 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
@@ -507,7 +507,8 @@ void NVPTXPassConfig::addMachineSSAOptimization() {
bool NVPTXTargetMachine::isCompatibleDataLayout(
const DataLayout &Candidate) const {
- //XXX: Should we enforce that the Candidate DataLayout has the same address
space for allocas?
+ // XXX: Should we enforce that the Candidate DataLayout has the same address
+ // space for allocas?
if (DL == Candidate)
return true;
@@ -520,9 +521,10 @@ bool NVPTXTargetMachine::isCompatibleDataLayout(
return NewDL == Candidate;
}
-unsigned NVPTXTargetMachine::getAddressSpaceForPseudoSourceKind(unsigned Kind)
const {
+unsigned
+NVPTXTargetMachine::getAddressSpaceForPseudoSourceKind(unsigned Kind) const {
if (Kind == PseudoSourceValue::FixedStack) {
return ADDRESS_SPACE_LOCAL;
}
return CodeGenTargetMachineImpl::getAddressSpaceForPseudoSourceKind(Kind);
-}
+}
\ No newline at end of file
>From 1c642db0d6fae1f2216ecdca71e662e623fa6bc4 Mon Sep 17 00:00:00 2001
From: Theodoros Theodoridis <ttheodori...@nvidia.com>
Date: Fri, 22 Aug 2025 10:19:18 +0000
Subject: [PATCH 3/7] Function -> Module Pass
---
llvm/lib/Target/NVPTX/NVPTX.h | 2 +-
llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp | 28 ++++++++++++----------
2 files changed, 16 insertions(+), 14 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTX.h b/llvm/lib/Target/NVPTX/NVPTX.h
index 77a0e03d4075a..b91af4d5a6305 100644
--- a/llvm/lib/Target/NVPTX/NVPTX.h
+++ b/llvm/lib/Target/NVPTX/NVPTX.h
@@ -48,7 +48,7 @@ MachineFunctionPass *createNVPTXPrologEpilogPass();
MachineFunctionPass *createNVPTXReplaceImageHandlesPass();
FunctionPass *createNVPTXImageOptimizerPass();
FunctionPass *createNVPTXLowerArgsPass();
-FunctionPass *createNVPTXLowerAllocaPass();
+ModulePass *createNVPTXLowerAllocaPass();
FunctionPass *createNVPTXLowerUnreachablePass(bool TrapUnreachable,
bool NoTrapAfterNoreturn);
FunctionPass *createNVPTXTagInvariantLoadsPass();
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
index 4477be9b65004..100d82219512e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
@@ -39,13 +39,14 @@ using namespace llvm;
using namespace NVPTXAS;
namespace {
-class NVPTXLowerAlloca : public FunctionPass {
- bool runOnFunction(Function &F) override;
- bool doInitialization(Module &M) override;
+class NVPTXLowerAlloca : public ModulePass {
+ bool changeDataLayout(Module &M);
+ bool lowerFunctionAllocas(Function &F);
public:
- static char ID; // Pass identification, replacement for typeid
- NVPTXLowerAlloca() : FunctionPass(ID) {}
+ static char ID;
+ NVPTXLowerAlloca() : ModulePass(ID) {}
+ bool runOnModule(Module &M) override;
StringRef getPassName() const override {
return "convert address space of alloca'ed memory to local";
}
@@ -57,13 +58,14 @@ char NVPTXLowerAlloca::ID = 1;
INITIALIZE_PASS(NVPTXLowerAlloca, "nvptx-lower-alloca", "Lower Alloca", false,
false)
-//
=============================================================================
-// Main function for this pass.
-//
=============================================================================
-bool NVPTXLowerAlloca::runOnFunction(Function &F) {
- if (skipFunction(F))
- return false;
+bool NVPTXLowerAlloca::runOnModule(Module &M) {
+ bool Changed = changeDataLayout(M);
+ for (auto &F : M)
+ Changed |= lowerFunctionAllocas(F);
+ return Changed;
+}
+bool NVPTXLowerAlloca::lowerFunctionAllocas(Function &F) {
SmallVector<AllocaInst *, 16> Allocas;
for (auto &BB : F)
for (auto &I : BB)
@@ -103,7 +105,7 @@ bool NVPTXLowerAlloca::runOnFunction(Function &F) {
return true;
}
-bool NVPTXLowerAlloca::doInitialization(Module &M) {
+bool NVPTXLowerAlloca::changeDataLayout(Module &M) {
const auto &DL = M.getDataLayout();
if (DL.getAllocaAddrSpace() == ADDRESS_SPACE_LOCAL)
return false;
@@ -115,6 +117,6 @@ bool NVPTXLowerAlloca::doInitialization(Module &M) {
return true;
}
-FunctionPass *llvm::createNVPTXLowerAllocaPass() {
+ModulePass *llvm::createNVPTXLowerAllocaPass() {
return new NVPTXLowerAlloca();
}
>From dd554137b792a7c70e83adc5918241877b6d5180 Mon Sep 17 00:00:00 2001
From: Theodoros Theodoridis <ttheodori...@nvidia.com>
Date: Fri, 22 Aug 2025 10:23:16 +0000
Subject: [PATCH 4/7] assert -> fatal_error
---
llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp | 6 ++++--
1 file changed, 4 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
index 100d82219512e..f4204294fe37e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXLowerAlloca.cpp
@@ -33,6 +33,7 @@
#include "llvm/IR/Module.h"
#include "llvm/IR/Type.h"
#include "llvm/Pass.h"
+#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/NVPTXAddrSpace.h"
using namespace llvm;
@@ -44,7 +45,7 @@ class NVPTXLowerAlloca : public ModulePass {
bool lowerFunctionAllocas(Function &F);
public:
- static char ID;
+ static char ID;
NVPTXLowerAlloca() : ModulePass(ID) {}
bool runOnModule(Module &M) override;
StringRef getPassName() const override {
@@ -112,7 +113,8 @@ bool NVPTXLowerAlloca::changeDataLayout(Module &M) {
auto DLStr = DL.getStringRepresentation();
auto AddrSpaceStr = "A" + std::to_string(ADDRESS_SPACE_LOCAL);
- assert(!StringRef(DLStr).contains("A") && "DataLayout should not contain A");
+ if (StringRef(DLStr).contains("A"))
+ report_fatal_error("DataLayout should not contain A");
M.setDataLayout(DLStr.empty() ? AddrSpaceStr : DLStr + "-" + AddrSpaceStr);
return true;
}
>From 94ee8c629b92a1a5a7e8d1a260608064b1192c70 Mon Sep 17 00:00:00 2001
From: Theodoros Theodoridis <ttheodori...@nvidia.com>
Date: Fri, 22 Aug 2025 10:37:30 +0000
Subject: [PATCH 5/7] Remove SP
---
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 12 +++++-------
llvm/lib/Target/NVPTX/NVPTXFrameLowering.cpp | 16 +++-------------
llvm/test/CodeGen/NVPTX/f32x2-instructions.ll | 1 -
llvm/test/CodeGen/NVPTX/indirect_byval.ll | 2 --
llvm/test/CodeGen/NVPTX/local-stack-frame.ll | 8 --------
.../CodeGen/NVPTX/lower-args-gridconstant.ll | 1 -
llvm/test/CodeGen/NVPTX/lower-byval-args.ll | 9 ---------
llvm/test/CodeGen/NVPTX/variadics-backend.ll | 5 -----
8 files changed, 8 insertions(+), 46 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index 8528cb41d9244..9ceabcbdd9c88 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -80,6 +80,7 @@
#include "llvm/Support/Compiler.h"
#include "llvm/Support/Endian.h"
#include "llvm/Support/ErrorHandling.h"
+#include "llvm/Support/NVPTXAddrSpace.h"
#include "llvm/Support/NativeFormatting.h"
#include "llvm/Support/raw_ostream.h"
#include "llvm/Target/TargetLoweringObjectFile.h"
@@ -1485,13 +1486,10 @@ void
NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters(
O << "\t.local .align " << MFI.getMaxAlign().value() << " .b8 \t"
<< DEPOTNAME << getFunctionNumber() << "[" << NumBytes << "];\n";
if (static_cast<const NVPTXTargetMachine &>(MF.getTarget())
- .getPointerSize(ADDRESS_SPACE_LOCAL) == 8) {
- O << "\t.reg .b64 \t%SP;\n"
- << "\t.reg .b64 \t%SPL;\n";
- } else {
- O << "\t.reg .b32 \t%SP;\n"
- << "\t.reg .b32 \t%SPL;\n";
- }
+ .getPointerSize(ADDRESS_SPACE_LOCAL) == 8)
+ O << "\t.reg .b64 \t%SPL;\n";
+ else
+ O << "\t.reg .b32 \t%SPL;\n";
}
// Go through all virtual registers to establish the mapping between the
diff --git a/llvm/lib/Target/NVPTX/NVPTXFrameLowering.cpp
b/llvm/lib/Target/NVPTX/NVPTXFrameLowering.cpp
index 78f18a93b869b..6e4817f09db6c 100644
--- a/llvm/lib/Target/NVPTX/NVPTXFrameLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXFrameLowering.cpp
@@ -46,21 +46,11 @@ void NVPTXFrameLowering::emitPrologue(MachineFunction &MF,
// Emits
// mov %SPL, %depot;
- // cvta.local %SP, %SPL;
// for local address accesses in MF.
- bool Is64Bit = static_cast<const NVPTXTargetMachine &>(MF.getTarget())
- .getPointerSize(NVPTXAS::ADDRESS_SPACE_LOCAL) == 8;
- unsigned CvtaLocalOpcode =
- (Is64Bit ? NVPTX::cvta_local_64 : NVPTX::cvta_local);
+ bool IsLocal64Bit = static_cast<const NVPTXTargetMachine &>(MF.getTarget())
+ .getPointerSize(NVPTXAS::ADDRESS_SPACE_LOCAL) == 8;
unsigned MovDepotOpcode =
- (Is64Bit ? NVPTX::MOV_DEPOT_ADDR_64 : NVPTX::MOV_DEPOT_ADDR);
- if (!MR.use_empty(NRI->getFrameRegister(MF))) {
- // If %SP is not used, do not bother emitting "cvta.local %SP, %SPL".
- MBBI = BuildMI(MBB, MBBI, dl,
- MF.getSubtarget().getInstrInfo()->get(CvtaLocalOpcode),
- NRI->getFrameRegister(MF))
- .addReg(NRI->getFrameLocalRegister(MF));
- }
+ (IsLocal64Bit ? NVPTX::MOV_DEPOT_ADDR_64 : NVPTX::MOV_DEPOT_ADDR);
if (!MR.use_empty(NRI->getFrameLocalRegister(MF))) {
BuildMI(MBB, MBBI, dl,
MF.getSubtarget().getInstrInfo()->get(MovDepotOpcode),
diff --git a/llvm/test/CodeGen/NVPTX/f32x2-instructions.ll
b/llvm/test/CodeGen/NVPTX/f32x2-instructions.ll
index cf00c2c9eaa3f..9b870c731cd0d 100644
--- a/llvm/test/CodeGen/NVPTX/f32x2-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/f32x2-instructions.ll
@@ -79,7 +79,6 @@ define float @test_extract_i(<2 x float> %a, i64 %idx) #0 {
; CHECK-NOF32X2-LABEL: test_extract_i(
; CHECK-NOF32X2: {
; CHECK-NOF32X2-NEXT: .local .align 8 .b8 __local_depot3[8];
-; CHECK-NOF32X2-NEXT: .reg .b64 %SP;
; CHECK-NOF32X2-NEXT: .reg .b64 %SPL;
; CHECK-NOF32X2-NEXT: .reg .b32 %r<4>;
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<6>;
diff --git a/llvm/test/CodeGen/NVPTX/indirect_byval.ll
b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
index 39813efef9b64..a27706854222d 100644
--- a/llvm/test/CodeGen/NVPTX/indirect_byval.ll
+++ b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
@@ -13,7 +13,6 @@ define internal i32 @foo() {
; CHECK-LABEL: foo(
; CHECK: {
; CHECK-NEXT: .local .align 1 .b8 __local_depot0[2];
-; CHECK-NEXT: .reg .b64 %SP;
; CHECK-NEXT: .reg .b64 %SPL;
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
@@ -50,7 +49,6 @@ define internal i32 @bar() {
; CHECK: // @bar
; CHECK-NEXT: {
; CHECK-NEXT: .local .align 8 .b8 __local_depot1[16];
-; CHECK-NEXT: .reg .b64 %SP;
; CHECK-NEXT: .reg .b64 %SPL;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-NEXT: .reg .b64 %rd<5>;
diff --git a/llvm/test/CodeGen/NVPTX/local-stack-frame.ll
b/llvm/test/CodeGen/NVPTX/local-stack-frame.ll
index 3e6f5e37aa85f..35bfde0dc2202 100644
--- a/llvm/test/CodeGen/NVPTX/local-stack-frame.ll
+++ b/llvm/test/CodeGen/NVPTX/local-stack-frame.ll
@@ -10,7 +10,6 @@ define void @foo(i32 %a) {
; PTX32-LABEL: foo(
; PTX32: {
; PTX32-NEXT: .local .align 4 .b8 __local_depot0[4];
-; PTX32-NEXT: .reg .b32 %SP;
; PTX32-NEXT: .reg .b32 %SPL;
; PTX32-NEXT: .reg .b32 %r<2>;
; PTX32-EMPTY:
@@ -23,7 +22,6 @@ define void @foo(i32 %a) {
; PTX64-LABEL: foo(
; PTX64: {
; PTX64-NEXT: .local .align 4 .b8 __local_depot0[4];
-; PTX64-NEXT: .reg .b64 %SP;
; PTX64-NEXT: .reg .b64 %SPL;
; PTX64-NEXT: .reg .b32 %r<2>;
; PTX64-EMPTY:
@@ -41,7 +39,6 @@ define ptx_kernel void @foo2(i32 %a) {
; PTX32-LABEL: foo2(
; PTX32: {
; PTX32-NEXT: .local .align 4 .b8 __local_depot1[4];
-; PTX32-NEXT: .reg .b32 %SP;
; PTX32-NEXT: .reg .b32 %SPL;
; PTX32-NEXT: .reg .b32 %r<4>;
; PTX32-EMPTY:
@@ -61,7 +58,6 @@ define ptx_kernel void @foo2(i32 %a) {
; PTX64-LABEL: foo2(
; PTX64: {
; PTX64-NEXT: .local .align 4 .b8 __local_depot1[4];
-; PTX64-NEXT: .reg .b64 %SP;
; PTX64-NEXT: .reg .b64 %SPL;
; PTX64-NEXT: .reg .b32 %r<2>;
; PTX64-NEXT: .reg .b64 %rd<3>;
@@ -90,7 +86,6 @@ define void @foo3(i32 %a) {
; PTX32-LABEL: foo3(
; PTX32: {
; PTX32-NEXT: .local .align 4 .b8 __local_depot2[12];
-; PTX32-NEXT: .reg .b32 %SP;
; PTX32-NEXT: .reg .b32 %SPL;
; PTX32-NEXT: .reg .b32 %r<5>;
; PTX32-EMPTY:
@@ -106,7 +101,6 @@ define void @foo3(i32 %a) {
; PTX64-LABEL: foo3(
; PTX64: {
; PTX64-NEXT: .local .align 4 .b8 __local_depot2[12];
-; PTX64-NEXT: .reg .b64 %SP;
; PTX64-NEXT: .reg .b64 %SPL;
; PTX64-NEXT: .reg .b32 %r<2>;
; PTX64-NEXT: .reg .b64 %rd<3>;
@@ -128,7 +122,6 @@ define void @foo4() {
; PTX32-LABEL: foo4(
; PTX32: {
; PTX32-NEXT: .local .align 4 .b8 __local_depot3[8];
-; PTX32-NEXT: .reg .b32 %SP;
; PTX32-NEXT: .reg .b32 %SPL;
; PTX32-NEXT: .reg .b32 %r<5>;
; PTX32-EMPTY:
@@ -155,7 +148,6 @@ define void @foo4() {
; PTX64-LABEL: foo4(
; PTX64: {
; PTX64-NEXT: .local .align 4 .b8 __local_depot3[8];
-; PTX64-NEXT: .reg .b64 %SP;
; PTX64-NEXT: .reg .b64 %SPL;
; PTX64-NEXT: .reg .b64 %rd<5>;
; PTX64-EMPTY:
diff --git a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
index f4a24ab41ba17..2cfca3f6b3174 100644
--- a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
@@ -149,7 +149,6 @@ define ptx_kernel void @multiple_grid_const_escape(ptr
byval(%struct.s) align 4
; PTX-LABEL: multiple_grid_const_escape(
; PTX: {
; PTX-NEXT: .local .align 4 .b8 __local_depot4[4];
-; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
; PTX-NEXT: .reg .b32 %r<2>;
; PTX-NEXT: .reg .b64 %rd<8>;
diff --git a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
index d8a9530e931a6..246eaf8a55538 100644
--- a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll
@@ -136,7 +136,6 @@ define dso_local ptx_kernel void @escape_ptr(ptr nocapture
noundef readnone %out
; PTX-LABEL: escape_ptr(
; PTX: {
; PTX-NEXT: .local .align 8 .b8 __local_depot2[8];
-; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
; PTX-NEXT: .reg .b64 %rd<7>;
; PTX-EMPTY:
@@ -175,7 +174,6 @@ define dso_local ptx_kernel void @escape_ptr_gep(ptr
nocapture noundef readnone
; PTX-LABEL: escape_ptr_gep(
; PTX: {
; PTX-NEXT: .local .align 8 .b8 __local_depot3[8];
-; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
; PTX-NEXT: .reg .b64 %rd<8>;
; PTX-EMPTY:
@@ -215,7 +213,6 @@ define dso_local ptx_kernel void @escape_ptr_store(ptr
nocapture noundef writeon
; PTX-LABEL: escape_ptr_store(
; PTX: {
; PTX-NEXT: .local .align 8 .b8 __local_depot4[8];
-; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
; PTX-NEXT: .reg .b64 %rd<9>;
; PTX-EMPTY:
@@ -252,7 +249,6 @@ 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 8 .b8 __local_depot5[8];
-; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
; PTX-NEXT: .reg .b64 %rd<10>;
; PTX-EMPTY:
@@ -291,7 +287,6 @@ define dso_local ptx_kernel void @escape_ptrtoint(ptr
nocapture noundef writeonl
; PTX-LABEL: escape_ptrtoint(
; PTX: {
; PTX-NEXT: .local .align 8 .b8 __local_depot6[8];
-; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
; PTX-NEXT: .reg .b64 %rd<9>;
; PTX-EMPTY:
@@ -448,7 +443,6 @@ define dso_local ptx_kernel void @memcpy_to_param(ptr
nocapture noundef readonly
; PTX-LABEL: memcpy_to_param(
; PTX: {
; PTX-NEXT: .local .align 8 .b8 __local_depot9[8];
-; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
; PTX-NEXT: .reg .b16 %rs<17>;
; PTX-NEXT: .reg .b64 %rd<8>;
@@ -628,7 +622,6 @@ define ptx_kernel void @test_select_write(ptr byval(i32)
align 4 %input1, ptr by
; PTX-LABEL: test_select_write(
; PTX: {
; PTX-NEXT: .local .align 4 .b8 __local_depot12[8];
-; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
; PTX-NEXT: .reg .pred %p<2>;
; PTX-NEXT: .reg .b16 %rs<3>;
@@ -806,7 +799,6 @@ define ptx_kernel void @test_phi_write(ptr byval(%struct.S)
align 4 %input1, ptr
; PTX-LABEL: test_phi_write(
; PTX: {
; PTX-NEXT: .local .align 4 .b8 __local_depot14[8];
-; PTX-NEXT: .reg .b64 %SP;
; PTX-NEXT: .reg .b64 %SPL;
; PTX-NEXT: .reg .pred %p<2>;
; PTX-NEXT: .reg .b16 %rs<3>;
@@ -858,7 +850,6 @@ define ptx_kernel void @test_forward_byval_arg(ptr
byval(i32) align 4 %input) {
; PTX-LABEL: test_forward_byval_arg(
; PTX: {
; 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-EMPTY:
diff --git a/llvm/test/CodeGen/NVPTX/variadics-backend.ll
b/llvm/test/CodeGen/NVPTX/variadics-backend.ll
index 7fe0dda2bf94e..e17d7d0b28269 100644
--- a/llvm/test/CodeGen/NVPTX/variadics-backend.ll
+++ b/llvm/test/CodeGen/NVPTX/variadics-backend.ll
@@ -102,7 +102,6 @@ define dso_local i32 @foo() {
; CHECK-PTX-LABEL: foo(
; CHECK-PTX: {
; CHECK-PTX-NEXT: .local .align 8 .b8 __local_depot1[40];
-; CHECK-PTX-NEXT: .reg .b64 %SP;
; CHECK-PTX-NEXT: .reg .b64 %SPL;
; CHECK-PTX-NEXT: .reg .b32 %r<2>;
; CHECK-PTX-NEXT: .reg .b64 %rd<3>;
@@ -139,7 +138,6 @@ define dso_local i32 @variadics2(i32 noundef %first, ...) {
; CHECK-PTX-LABEL: variadics2(
; CHECK-PTX: {
; 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<6>;
; CHECK-PTX-NEXT: .reg .b32 %r<6>;
@@ -199,7 +197,6 @@ define dso_local i32 @bar() {
; CHECK-PTX-LABEL: bar(
; CHECK-PTX: {
; 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<6>;
; CHECK-PTX-NEXT: .reg .b32 %r<2>;
@@ -280,7 +277,6 @@ define dso_local i32 @baz() {
; CHECK-PTX-LABEL: baz(
; CHECK-PTX: {
; CHECK-PTX-NEXT: .local .align 16 .b8 __local_depot5[16];
-; CHECK-PTX-NEXT: .reg .b64 %SP;
; CHECK-PTX-NEXT: .reg .b64 %SPL;
; CHECK-PTX-NEXT: .reg .b32 %r<2>;
; CHECK-PTX-NEXT: .reg .b64 %rd<3>;
@@ -346,7 +342,6 @@ define dso_local void @qux() {
; CHECK-PTX-LABEL: qux(
; CHECK-PTX: {
; 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<5>;
; CHECK-PTX-EMPTY:
>From b37221cf6d3e53d914143106fb9e968717bd671f Mon Sep 17 00:00:00 2001
From: Theodoros Theodoridis <ttheodori...@nvidia.com>
Date: Fri, 22 Aug 2025 10:40:02 +0000
Subject: [PATCH 6/7] Fix pointer size query
---
llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
index 04eb095f31180..709a5832bb784 100644
--- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
@@ -128,7 +128,7 @@ bool
NVPTXRegisterInfo::eliminateFrameIndex(MachineBasicBlock::iterator II,
Register NVPTXRegisterInfo::getFrameRegister(const MachineFunction &MF) const {
const NVPTXTargetMachine &TM =
static_cast<const NVPTXTargetMachine &>(MF.getTarget());
- return TM.getPointerSize(NVPTXAS::ADDRESS_SPACE_LOCAL) == 8
+ return TM.getPointerSize(NVPTXAS::ADDRESS_SPACE_GENERIC) == 8
? NVPTX::VRFrame64
: NVPTX::VRFrame32;
}
>From 4a15fd4f4856fb98cf040623d90c6e8e90c427e9 Mon Sep 17 00:00:00 2001
From: Theodoros Theodoridis <ttheodori...@nvidia.com>
Date: Fri, 22 Aug 2025 12:30:02 +0000
Subject: [PATCH 7/7] [clang] modify backend datalayout check
---
clang/lib/CodeGen/BackendUtil.cpp | 13 +++++++------
1 file changed, 7 insertions(+), 6 deletions(-)
diff --git a/clang/lib/CodeGen/BackendUtil.cpp
b/clang/lib/CodeGen/BackendUtil.cpp
index 0b8b824fbcd5a..8198f9c58340c 100644
--- a/clang/lib/CodeGen/BackendUtil.cpp
+++ b/clang/lib/CodeGen/BackendUtil.cpp
@@ -1444,15 +1444,16 @@ void clang::emitBackendOutput(CompilerInstance &CI,
CodeGenOptions &CGOpts,
// Verify clang's TargetInfo DataLayout against the LLVM TargetMachine's
// DataLayout.
- if (AsmHelper.TM) {
- std::string DLDesc = M->getDataLayout().getStringRepresentation();
- if (DLDesc != TDesc) {
+ if (AsmHelper.TM)
+ if (!AsmHelper.TM->isCompatibleDataLayout(M->getDataLayout()) ||
+ !AsmHelper.TM->isCompatibleDataLayout(DataLayout(TDesc))) {
+ std::string DLDesc = M->getDataLayout().getStringRepresentation();
unsigned DiagID = Diags.getCustomDiagID(
- DiagnosticsEngine::Error, "backend data layout '%0' does not match "
- "expected target description '%1'");
+ DiagnosticsEngine::Error,
+ "backend data layout '%0' is not compatible with "
+ "expected target description '%1'");
Diags.Report(DiagID) << DLDesc << TDesc;
}
- }
}
// With -fembed-bitcode, save a copy of the llvm IR as data in the
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
