[llvm] 94c1a8b - [NVPTX] Fold symbol addresses into memory operands (#202379)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Jun 25 08:22:37 PDT 2026
Author: Tim Besard
Date: 2026-06-25T08:22:31-07:00
New Revision: 94c1a8bc3e6b790ccc5723803129b3b956f11d13
URL: https://github.com/llvm/llvm-project/commit/94c1a8bc3e6b790ccc5723803129b3b956f11d13
DIFF: https://github.com/llvm/llvm-project/commit/94c1a8bc3e6b790ccc5723803129b3b956f11d13.diff
LOG: [NVPTX] Fold symbol addresses into memory operands (#202379)
SelectionDAG can fold a symbol address (a kernel parameter, global
variable, or external symbol) directly into a memory instruction's
address operand, but only within a single basic block. When the address
crosses a block boundary, ISel materializes it with `MOV_B{32,64}_sym`
and the memory instruction becomes register-relative:
```ptx
mov.b64 %rd1, kernel_param_0;
ld.param.b64 %rd2, [%rd1];
ld.param.b64 %rd3, [%rd1+8];
```
instead of:
```
ld.param.b64 %rd2, [kernel_param_0];
ld.param.b64 %rd3, [kernel_param_0+8];
```
This patch adds NVPTXAddressFolder, a pre-regalloc pass that looks for
loads and stores whose address operand is defined by `MOV_B{32,64}_sym`,
then folds the symbol back into the memory operand. The mov is erased
once it has no remaining uses; if the address also feeds arithmetic or
escapes, it is kept.
To make this generic over NVPTX memory instructions, the patch enables
named operand tables for NVPTX instructions and uses them to find `addr`
and `addsp` operands instead of hardcoding opcode-specific operand
indices.
This was motivated by a CUDA.jl performance regression where `byval`
kernel parameters stopped being pre-lowered into simple loads and
exposed the missing cross-block fold in the backend.
Disclaimer: LLMs (GPT 5.5, Opus 4.8) were used to develop this PR.
---------
Co-authored-by: Claude Opus 4.8 (1M context) <noreply at anthropic.com>
Added:
llvm/lib/Target/NVPTX/NVPTXAddressFolder.cpp
llvm/test/CodeGen/NVPTX/address-folder.ll
llvm/test/CodeGen/NVPTX/address-folder.mir
Modified:
llvm/lib/Target/NVPTX/CMakeLists.txt
llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp
llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.h
llvm/lib/Target/NVPTX/NVPTX.h
llvm/lib/Target/NVPTX/NVPTXInstrFormats.td
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp
Removed:
################################################################################
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
+...
More information about the llvm-commits
mailing list