[llvm] [NVPTX] Fold symbol addresses into memory operands (reland) (PR #205983)
Alex MacLean via llvm-commits
llvm-commits at lists.llvm.org
Fri Jun 26 08:04:10 PDT 2026
https://github.com/AlexMaclean updated https://github.com/llvm/llvm-project/pull/205983
>From a331ba799c2e03915d5a78a30e3ba073259bcdeb Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Thu, 25 Jun 2026 22:40:21 +0200
Subject: [PATCH 1/2] Reapply "[NVPTX] Fold symbol addresses into memory
operands" (#202379)
This reverts commit b7c56b5b0a54a6c966a74dfe311fcddcffade521, reapplying
the original change from #202379.
The original commit was reverted because CodeGen/NVPTX/address-folder.ll
failed on the buildbots. The failure was not in the pass itself but in the
test's hard-coded -O0 register counts, which had gone stale: the test was
generated and last CI-tested (premerge) on a base that predated #204192
([NVPTX] Rewrite kernel signatures in param AS). #204192 landed between this
PR's last premerge run and its merge, and it changed -O0 byval lowering so
that one redundant cross-block copy is no longer emitted, lowering the b64
register count by one. The -O0 checks are regenerated against current main
in the following commit.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply at anthropic.com>
---
llvm/lib/Target/NVPTX/CMakeLists.txt | 1 +
.../NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp | 1 +
.../NVPTX/MCTargetDesc/NVPTXMCTargetDesc.h | 2 +
llvm/lib/Target/NVPTX/NVPTX.h | 3 +
llvm/lib/Target/NVPTX/NVPTXAddressFolder.cpp | 123 ++++++++++++++++++
llvm/lib/Target/NVPTX/NVPTXInstrFormats.td | 1 +
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 4 +-
llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp | 3 +
llvm/test/CodeGen/NVPTX/address-folder.ll | 113 ++++++++++++++++
llvm/test/CodeGen/NVPTX/address-folder.mir | 60 +++++++++
10 files changed, 309 insertions(+), 2 deletions(-)
create mode 100644 llvm/lib/Target/NVPTX/NVPTXAddressFolder.cpp
create mode 100644 llvm/test/CodeGen/NVPTX/address-folder.ll
create mode 100644 llvm/test/CodeGen/NVPTX/address-folder.mir
diff --git a/llvm/lib/Target/NVPTX/CMakeLists.txt b/llvm/lib/Target/NVPTX/CMakeLists.txt
index e0e2b1a40bffc..d2b5083a5f83f 100644
--- a/llvm/lib/Target/NVPTX/CMakeLists.txt
+++ b/llvm/lib/Target/NVPTX/CMakeLists.txt
@@ -12,6 +12,7 @@ tablegen(LLVM NVPTXGenSubtargetInfo.inc -gen-subtarget)
add_public_tablegen_target(NVPTXCommonTableGen)
set(NVPTXCodeGen_sources
+ NVPTXAddressFolder.cpp
NVPTXAliasAnalysis.cpp
NVPTXAllocaHoisting.cpp
NVPTXAsmPrinter.cpp
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp
index 5933f74fcb25f..ef447e4ceed6b 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp
@@ -24,6 +24,7 @@
using namespace llvm;
#define GET_INSTRINFO_MC_DESC
+#define GET_INSTRINFO_NAMED_OPS
#define ENABLE_INSTR_PREDICATE_VERIFIER
#include "NVPTXGenInstrInfo.inc"
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.h b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.h
index 78f4e67455023..58083e1248bb0 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.h
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.h
@@ -13,6 +13,7 @@
#ifndef LLVM_LIB_TARGET_NVPTX_MCTARGETDESC_NVPTXMCTARGETDESC_H
#define LLVM_LIB_TARGET_NVPTX_MCTARGETDESC_NVPTXMCTARGETDESC_H
+#include "llvm/Support/Compiler.h"
#include <stdint.h>
// Defines symbolic names for PTX registers.
@@ -22,6 +23,7 @@
// Defines symbolic names for the PTX instructions.
#define GET_INSTRINFO_ENUM
#define GET_INSTRINFO_MC_HELPER_DECLS
+#define GET_INSTRINFO_OPERAND_ENUM
#include "NVPTXGenInstrInfo.inc"
#define GET_SUBTARGETINFO_ENUM
diff --git a/llvm/lib/Target/NVPTX/NVPTX.h b/llvm/lib/Target/NVPTX/NVPTX.h
index 08ccc143dfb8a..57ed9bb43aeea 100644
--- a/llvm/lib/Target/NVPTX/NVPTX.h
+++ b/llvm/lib/Target/NVPTX/NVPTX.h
@@ -59,6 +59,7 @@ FunctionPass *createNVPTXIRPeepholePass();
MachineFunctionPass *createNVPTXPeephole();
MachineFunctionPass *createNVPTXProxyRegErasurePass();
MachineFunctionPass *createNVPTXForwardParamsPass();
+MachineFunctionPass *createNVPTXAddressFolderPass();
void initializeNVVMReflectLegacyPassPass(PassRegistry &);
void initializeGenericToNVVMLegacyPassPass(PassRegistry &);
@@ -74,6 +75,7 @@ void initializeNVPTXLowerArgsLegacyPassPass(PassRegistry &);
void initializeNVPTXSetByValParamAlignLegacyPassPass(PassRegistry &);
void initializeNVPTXProxyRegErasurePass(PassRegistry &);
void initializeNVPTXForwardParamsPassPass(PassRegistry &);
+void initializeNVPTXAddressFolderPassPass(PassRegistry &);
void initializeNVVMIntrRangePass(PassRegistry &);
void initializeNVVMReflectPass(PassRegistry &);
void initializeNVPTXAAWrapperPassPass(PassRegistry &);
@@ -294,6 +296,7 @@ void initializeNVPTXDAGToDAGISelLegacyPass(PassRegistry &);
// Defines symbolic names for the NVPTX instructions.
#define GET_INSTRINFO_ENUM
#define GET_INSTRINFO_MC_HELPER_DECLS
+#define GET_INSTRINFO_OPERAND_ENUM
#include "NVPTXGenInstrInfo.inc"
#endif
diff --git a/llvm/lib/Target/NVPTX/NVPTXAddressFolder.cpp b/llvm/lib/Target/NVPTX/NVPTXAddressFolder.cpp
new file mode 100644
index 0000000000000..5299dbe686611
--- /dev/null
+++ b/llvm/lib/Target/NVPTX/NVPTXAddressFolder.cpp
@@ -0,0 +1,123 @@
+//===- NVPTXAddressFolder.cpp - Fold symbol addresses into memory ops -----===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// SelectionDAG folds a symbol address (a kernel parameter, global variable, or
+// external symbol) directly into the address operand of a memory access, but
+// only within a single basic block. When the address is carried across a block
+// boundary in a register it is materialized with a generic `mov`
+// (MOV_B{32,64}_sym) and the accesses become register-relative:
+//
+// mov.b64 %rd1, kernel_param_0;
+// ld.param.b64 %rd2, [%rd1];
+// ld.param.b64 %rd3, [%rd1+8];
+//
+// This pass folds the symbol back into those address operands, eliminating the
+// redundant address arithmetic (and the `mov` itself once no use remains):
+//
+// ld.param.b64 %rd2, [kernel_param_0];
+// ld.param.b64 %rd3, [kernel_param_0+8];
+//
+// Shared-memory accesses are left alone: `mov`s of shared symbols are
+// deliberately kept CSE-able rather than duplicated into their uses, as
+// rematerializing them has caused performance regressions before (see
+// MovSymInst in NVPTXInstrInfo.td).
+//
+//===----------------------------------------------------------------------===//
+
+#include "NVPTX.h"
+#include "llvm/CodeGen/MachineFunctionPass.h"
+#include "llvm/CodeGen/MachineInstr.h"
+#include "llvm/CodeGen/MachineOperand.h"
+#include "llvm/CodeGen/MachineRegisterInfo.h"
+
+using namespace llvm;
+
+// Try to fold the definition of \p Addr into \p MI's address operand. The
+// defining instruction is erased once it has no uses left; it is kept for any
+// remaining use, e.g. because the address also feeds arithmetic or escapes.
+static bool foldAddress(MachineInstr &MI, MachineOperand &Addr,
+ MachineRegisterInfo &MRI) {
+ assert(Addr.isReg() && "Expected an address register");
+ assert(Addr.getReg().isVirtual() && "Expected a virtual address register");
+
+ MachineInstr *Mov = MRI.getVRegDef(Addr.getReg());
+ if (!Mov || (Mov->getOpcode() != NVPTX::MOV_B32_sym &&
+ Mov->getOpcode() != NVPTX::MOV_B64_sym))
+ return false;
+
+ const MachineOperand &Sym = Mov->getOperand(1);
+ if (!Sym.isGlobal() && !Sym.isSymbol())
+ return false;
+
+ // The accessed address space must be known and must not be shared.
+ const int AddrSpaceIdx =
+ NVPTX::getNamedOperandIdx(MI.getOpcode(), NVPTX::OpName::addsp);
+ if (AddrSpaceIdx < 0)
+ return false;
+ const auto AddrSpace = MI.getOperand(AddrSpaceIdx).getImm();
+ if (AddrSpace == NVPTX::AddressSpace::Shared ||
+ AddrSpace == NVPTX::AddressSpace::SharedCluster)
+ return false;
+
+ if (Sym.isGlobal()) {
+ Addr.ChangeToGA(Sym.getGlobal(), Sym.getOffset(), Sym.getTargetFlags());
+ } else {
+ Addr.ChangeToES(Sym.getSymbolName(), Sym.getTargetFlags());
+ }
+
+ if (MRI.use_empty(Mov->getOperand(0).getReg()))
+ Mov->eraseFromParent();
+
+ return true;
+}
+
+static bool foldAddresses(MachineFunction &MF) {
+ MachineRegisterInfo &MRI = MF.getRegInfo();
+
+ bool Changed = false;
+ for (MachineBasicBlock &MBB : MF)
+ for (MachineInstr &MI : make_early_inc_range(MBB))
+ if (MI.mayLoadOrStore()) {
+ const int AddrIdx =
+ NVPTX::getNamedOperandIdx(MI.getOpcode(), NVPTX::OpName::addr);
+ if (AddrIdx >= 0 && MI.getOperand(AddrIdx).isReg())
+ Changed |= foldAddress(MI, MI.getOperand(AddrIdx), MRI);
+ }
+
+ return Changed;
+}
+
+/// ----------------------------------------------------------------------------
+/// Pass (Manager) Boilerplate
+/// ----------------------------------------------------------------------------
+
+namespace {
+struct NVPTXAddressFolderPass : public MachineFunctionPass {
+ static char ID;
+ NVPTXAddressFolderPass() : MachineFunctionPass(ID) {}
+
+ bool runOnMachineFunction(MachineFunction &MF) override;
+
+ void getAnalysisUsage(AnalysisUsage &AU) const override {
+ MachineFunctionPass::getAnalysisUsage(AU);
+ }
+};
+} // namespace
+
+char NVPTXAddressFolderPass::ID = 0;
+
+INITIALIZE_PASS(NVPTXAddressFolderPass, "nvptx-address-folder",
+ "NVPTX Address Folder", false, false)
+
+bool NVPTXAddressFolderPass::runOnMachineFunction(MachineFunction &MF) {
+ return foldAddresses(MF);
+}
+
+MachineFunctionPass *llvm::createNVPTXAddressFolderPass() {
+ return new NVPTXAddressFolderPass();
+}
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrFormats.td b/llvm/lib/Target/NVPTX/NVPTXInstrFormats.td
index 719be0300940e..361f819a340ae 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrFormats.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrFormats.td
@@ -22,6 +22,7 @@ class NVPTXInst<dag outs, dag ins, string asmstr, list<dag> pattern = []>
dag InOperandList = ins;
let AsmString = asmstr;
let Pattern = pattern;
+ let UseNamedOperandTable = true;
// TSFlagFields
bit IsLoad = false;
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 5f1c16dc7f0fb..c9db5c5cec025 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -2454,8 +2454,8 @@ let isTerminator=1 in {
let isBranch=1 in {
def CBranch : NVPTXInst<(outs),
- (ins B1:$a, brtarget:$target, BranchFlag:$not),
- "@${not}$a bra \t$target;",
+ (ins B1:$a, brtarget:$target, BranchFlag:$negate),
+ "@${negate}$a bra \t$target;",
[(brcond i1:$a, bb:$target)]>;
let isBarrier=1 in
diff --git a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
index 6054feeedf932..2bc2158fccae1 100644
--- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
@@ -119,6 +119,7 @@ extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializeNVPTXTarget() {
initializeNVPTXLowerAggrCopiesPass(PR);
initializeNVPTXProxyRegErasurePass(PR);
initializeNVPTXForwardParamsPassPass(PR);
+ initializeNVPTXAddressFolderPassPass(PR);
initializeNVPTXDAGToDAGISelLegacyPass(PR);
initializeNVPTXAAWrapperPassPass(PR);
initializeNVPTXExternalAAWrapperPass(PR);
@@ -421,6 +422,8 @@ bool NVPTXPassConfig::addInstSelector() {
void NVPTXPassConfig::addPreRegAlloc() {
addPass(createNVPTXForwardParamsPass());
+ if (getOptLevel() != CodeGenOptLevel::None)
+ addPass(createNVPTXAddressFolderPass());
// Remove Proxy Register pseudo instructions used to keep `callseq_end` alive.
addPass(createNVPTXProxyRegErasurePass());
}
diff --git a/llvm/test/CodeGen/NVPTX/address-folder.ll b/llvm/test/CodeGen/NVPTX/address-folder.ll
new file mode 100644
index 0000000000000..239f21594265d
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/address-folder.ll
@@ -0,0 +1,113 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s | FileCheck %s
+; RUN: llc -O0 < %s | FileCheck %s --check-prefix=O0
+
+target triple = "nvptx64-nvidia-cuda"
+
+; A byval kernel parameter loaded at a constant offset from a block other than
+; the one that defines its address must still be read directly from parameter
+; space (ld.param [param+off]), not via a materialized base register.
+define ptx_kernel void @test_kernel_cross_block(ptr byval([2 x i64]) align 8 %r, i64 %n, ptr addrspace(1) %out) {
+; CHECK-LABEL: test_kernel_cross_block(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b64 %rd<6>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd2, [test_kernel_cross_block_param_1];
+; CHECK-NEXT: setp.lt.s64 %p1, %rd2, 1;
+; CHECK-NEXT: @%p1 bra $L__BB0_2;
+; CHECK-NEXT: // %bb.1: // %body
+; CHECK-NEXT: ld.param.b64 %rd1, [test_kernel_cross_block_param_2];
+; CHECK-NEXT: ld.param.b64 %rd3, [test_kernel_cross_block_param_0];
+; CHECK-NEXT: ld.param.b64 %rd4, [test_kernel_cross_block_param_0+8];
+; CHECK-NEXT: mul.lo.s64 %rd5, %rd3, %rd4;
+; CHECK-NEXT: st.global.b64 [%rd1], %rd5;
+; CHECK-NEXT: $L__BB0_2: // %done
+; CHECK-NEXT: ret;
+;
+; O0-LABEL: test_kernel_cross_block(
+; O0: {
+; O0-NEXT: .reg .pred %p<2>;
+; O0-NEXT: .reg .b64 %rd<8>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd4, [test_kernel_cross_block_param_2];
+; O0-NEXT: ld.param.b64 %rd3, [test_kernel_cross_block_param_1];
+; O0-NEXT: mov.b64 %rd1, test_kernel_cross_block_param_0;
+; O0-NEXT: mov.b64 %rd2, %rd1;
+; O0-NEXT: setp.lt.s64 %p1, %rd3, 1;
+; O0-NEXT: @%p1 bra $L__BB0_2;
+; O0-NEXT: bra.uni $L__BB0_1;
+; O0-NEXT: $L__BB0_1: // %body
+; O0-NEXT: ld.param.b64 %rd5, [%rd1];
+; O0-NEXT: ld.param.b64 %rd6, [%rd1+8];
+; O0-NEXT: mul.lo.s64 %rd7, %rd5, %rd6;
+; O0-NEXT: st.global.b64 [%rd4], %rd7;
+; O0-NEXT: bra.uni $L__BB0_2;
+; O0-NEXT: $L__BB0_2: // %done
+; O0-NEXT: ret;
+ %c = icmp sgt i64 %n, 0
+ br i1 %c, label %body, label %done
+body:
+ %v0 = load i64, ptr %r, align 8
+ %p1 = getelementptr inbounds i8, ptr %r, i64 8
+ %v1 = load i64, ptr %p1, align 8
+ %m = mul i64 %v0, %v1
+ store i64 %m, ptr addrspace(1) %out, align 8
+ br label %done
+done:
+ ret void
+}
+
+; When the address is also used for arithmetic (here a variable offset) it is
+; not just loaded, so the `mov` must be kept rather than folded into the load.
+define ptx_kernel void @test_kernel_var_offset(ptr byval([8 x i64]) align 8 %r, i64 %n, ptr addrspace(1) %out) {
+; CHECK-LABEL: test_kernel_var_offset(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b64 %rd<6>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd4, [test_kernel_var_offset_param_1];
+; CHECK-NEXT: setp.lt.s64 %p1, %rd4, 1;
+; CHECK-NEXT: @%p1 bra $L__BB1_2;
+; CHECK-NEXT: // %bb.1: // %body
+; CHECK-NEXT: ld.param.b64 %rd2, [test_kernel_var_offset_param_2];
+; CHECK-NEXT: mov.b64 %rd3, test_kernel_var_offset_param_0;
+; CHECK-NEXT: add.s64 %rd1, %rd3, %rd4;
+; CHECK-NEXT: ld.param.b64 %rd5, [%rd1];
+; CHECK-NEXT: st.global.b64 [%rd2], %rd5;
+; CHECK-NEXT: $L__BB1_2: // %done
+; CHECK-NEXT: ret;
+;
+; O0-LABEL: test_kernel_var_offset(
+; O0: {
+; O0-NEXT: .reg .pred %p<2>;
+; O0-NEXT: .reg .b64 %rd<7>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd4, [test_kernel_var_offset_param_2];
+; O0-NEXT: ld.param.b64 %rd3, [test_kernel_var_offset_param_1];
+; O0-NEXT: mov.b64 %rd1, test_kernel_var_offset_param_0;
+; O0-NEXT: mov.b64 %rd2, %rd1;
+; O0-NEXT: setp.lt.s64 %p1, %rd3, 1;
+; O0-NEXT: @%p1 bra $L__BB1_2;
+; O0-NEXT: bra.uni $L__BB1_1;
+; O0-NEXT: $L__BB1_1: // %body
+; O0-NEXT: add.s64 %rd5, %rd1, %rd3;
+; O0-NEXT: ld.param.b64 %rd6, [%rd5];
+; O0-NEXT: st.global.b64 [%rd4], %rd6;
+; O0-NEXT: bra.uni $L__BB1_2;
+; O0-NEXT: $L__BB1_2: // %done
+; O0-NEXT: ret;
+ %c = icmp sgt i64 %n, 0
+ br i1 %c, label %body, label %done
+body:
+ %p = getelementptr inbounds i8, ptr %r, i64 %n
+ %v = load i64, ptr %p, align 8
+ store i64 %v, ptr addrspace(1) %out, align 8
+ br label %done
+done:
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/address-folder.mir b/llvm/test/CodeGen/NVPTX/address-folder.mir
new file mode 100644
index 0000000000000..ac50925579506
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/address-folder.mir
@@ -0,0 +1,60 @@
+# RUN: llc -mtriple=nvptx64-nvidia-cuda -run-pass=nvptx-address-folder -verify-machineinstrs -o - %s | FileCheck %s
+
+--- |
+ target triple = "nvptx64-nvidia-cuda"
+
+ @g = addrspace(1) global [4 x i64] zeroinitializer
+ @g_shared = addrspace(3) global [4 x i64] zeroinitializer
+
+ define void @fold_global() { ret void }
+ define void @skip_shared() { ret void }
+ define void @fold_partial() { ret void }
+...
+---
+# A global address materialized into a register is folded into the address
+# operand of loads and stores, and the dead `mov` is removed.
+# CHECK-LABEL: name: fold_global
+# CHECK-NOT: MOV_B64_sym
+# CHECK: %1:b64 = LD_i64 0, 0, 1, 3, 64, -1, @g, 0
+# CHECK: %2:b64 = LD_i64 0, 0, 1, 3, 64, -1, @g, 8
+# CHECK: ST_i64 %1, 0, 0, 1, 64, @g, 16
+name: fold_global
+tracksRegLiveness: true
+body: |
+ bb.0:
+ %0:b64 = MOV_B64_sym @g
+ %1:b64 = LD_i64 0, 0, 1, 3, 64, -1, %0, 0 :: (load (s64), addrspace 1)
+ %2:b64 = LD_i64 0, 0, 1, 3, 64, -1, %0, 8 :: (load (s64), addrspace 1)
+ ST_i64 %1, 0, 0, 1, 64, %0, 16 :: (store (s64), addrspace 1)
+ Return
+...
+---
+# Shared-memory accesses are not folded: the `mov` of the shared symbol must
+# stay CSE-able rather than be duplicated into every access.
+# CHECK-LABEL: name: skip_shared
+# CHECK: %0:b64 = MOV_B64_sym @g_shared
+# CHECK: %1:b64 = LD_i64 0, 0, 3, 3, 64, -1, %0, 0
+name: skip_shared
+tracksRegLiveness: true
+body: |
+ bb.0:
+ %0:b64 = MOV_B64_sym @g_shared
+ %1:b64 = LD_i64 0, 0, 3, 3, 64, -1, %0, 0 :: (load (s64), addrspace 3)
+ Return
+...
+---
+# When the address has non-memory uses too, memory uses are still folded but
+# the `mov` is kept for the remaining use.
+# CHECK-LABEL: name: fold_partial
+# CHECK: %0:b64 = MOV_B64_sym @g
+# CHECK: %1:b64 = LD_i64 0, 0, 1, 3, 64, -1, @g, 0
+# CHECK: %2:b64 = MULT64rr %0, %0
+name: fold_partial
+tracksRegLiveness: true
+body: |
+ bb.0:
+ %0:b64 = MOV_B64_sym @g
+ %1:b64 = LD_i64 0, 0, 1, 3, 64, -1, %0, 0 :: (load (s64), addrspace 1)
+ %2:b64 = MULT64rr %0, %0
+ Return
+...
>From ddf1b128a03833ec7ee7576555c0bd98553fb69a Mon Sep 17 00:00:00 2001
From: Tim Besard <tim.besard at gmail.com>
Date: Thu, 25 Jun 2026 22:53:18 +0200
Subject: [PATCH 2/2] [NVPTX] Regenerate address-folder.ll -O0 checks
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply at anthropic.com>
---
llvm/test/CodeGen/NVPTX/address-folder.ll | 32 +++++++++++------------
1 file changed, 15 insertions(+), 17 deletions(-)
diff --git a/llvm/test/CodeGen/NVPTX/address-folder.ll b/llvm/test/CodeGen/NVPTX/address-folder.ll
index 239f21594265d..813fd2c439651 100644
--- a/llvm/test/CodeGen/NVPTX/address-folder.ll
+++ b/llvm/test/CodeGen/NVPTX/address-folder.ll
@@ -29,21 +29,20 @@ define ptx_kernel void @test_kernel_cross_block(ptr byval([2 x i64]) align 8 %r,
; O0-LABEL: test_kernel_cross_block(
; O0: {
; O0-NEXT: .reg .pred %p<2>;
-; O0-NEXT: .reg .b64 %rd<8>;
+; O0-NEXT: .reg .b64 %rd<7>;
; O0-EMPTY:
; O0-NEXT: // %bb.0:
-; O0-NEXT: ld.param.b64 %rd4, [test_kernel_cross_block_param_2];
-; O0-NEXT: ld.param.b64 %rd3, [test_kernel_cross_block_param_1];
+; O0-NEXT: ld.param.b64 %rd3, [test_kernel_cross_block_param_2];
+; O0-NEXT: ld.param.b64 %rd2, [test_kernel_cross_block_param_1];
; O0-NEXT: mov.b64 %rd1, test_kernel_cross_block_param_0;
-; O0-NEXT: mov.b64 %rd2, %rd1;
-; O0-NEXT: setp.lt.s64 %p1, %rd3, 1;
+; O0-NEXT: setp.lt.s64 %p1, %rd2, 1;
; O0-NEXT: @%p1 bra $L__BB0_2;
; O0-NEXT: bra.uni $L__BB0_1;
; O0-NEXT: $L__BB0_1: // %body
-; O0-NEXT: ld.param.b64 %rd5, [%rd1];
-; O0-NEXT: ld.param.b64 %rd6, [%rd1+8];
-; O0-NEXT: mul.lo.s64 %rd7, %rd5, %rd6;
-; O0-NEXT: st.global.b64 [%rd4], %rd7;
+; O0-NEXT: ld.param.b64 %rd4, [%rd1];
+; O0-NEXT: ld.param.b64 %rd5, [%rd1+8];
+; O0-NEXT: mul.lo.s64 %rd6, %rd4, %rd5;
+; O0-NEXT: st.global.b64 [%rd3], %rd6;
; O0-NEXT: bra.uni $L__BB0_2;
; O0-NEXT: $L__BB0_2: // %done
; O0-NEXT: ret;
@@ -84,20 +83,19 @@ define ptx_kernel void @test_kernel_var_offset(ptr byval([8 x i64]) align 8 %r,
; O0-LABEL: test_kernel_var_offset(
; O0: {
; O0-NEXT: .reg .pred %p<2>;
-; O0-NEXT: .reg .b64 %rd<7>;
+; O0-NEXT: .reg .b64 %rd<6>;
; O0-EMPTY:
; O0-NEXT: // %bb.0:
-; O0-NEXT: ld.param.b64 %rd4, [test_kernel_var_offset_param_2];
-; O0-NEXT: ld.param.b64 %rd3, [test_kernel_var_offset_param_1];
+; O0-NEXT: ld.param.b64 %rd3, [test_kernel_var_offset_param_2];
+; O0-NEXT: ld.param.b64 %rd2, [test_kernel_var_offset_param_1];
; O0-NEXT: mov.b64 %rd1, test_kernel_var_offset_param_0;
-; O0-NEXT: mov.b64 %rd2, %rd1;
-; O0-NEXT: setp.lt.s64 %p1, %rd3, 1;
+; O0-NEXT: setp.lt.s64 %p1, %rd2, 1;
; O0-NEXT: @%p1 bra $L__BB1_2;
; O0-NEXT: bra.uni $L__BB1_1;
; O0-NEXT: $L__BB1_1: // %body
-; O0-NEXT: add.s64 %rd5, %rd1, %rd3;
-; O0-NEXT: ld.param.b64 %rd6, [%rd5];
-; O0-NEXT: st.global.b64 [%rd4], %rd6;
+; O0-NEXT: add.s64 %rd4, %rd1, %rd2;
+; O0-NEXT: ld.param.b64 %rd5, [%rd4];
+; O0-NEXT: st.global.b64 [%rd3], %rd5;
; O0-NEXT: bra.uni $L__BB1_2;
; O0-NEXT: $L__BB1_2: // %done
; O0-NEXT: ret;
More information about the llvm-commits
mailing list