Author: Krzysztof Parzyszek Date: 2026-06-25T11:20:14-05:00 New Revision: 1508332f203092dda1753eafc206af0d2ebbd0c2
URL: https://github.com/llvm/llvm-project/commit/1508332f203092dda1753eafc206af0d2ebbd0c2 DIFF: https://github.com/llvm/llvm-project/commit/1508332f203092dda1753eafc206af0d2ebbd0c2.diff LOG: Revert "[NVPTX] Fold symbol addresses into memory operands (#202379)" This reverts commit 94c1a8bc3e6b790ccc5723803129b3b956f11d13. Added: 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: llvm/lib/Target/NVPTX/NVPTXAddressFolder.cpp llvm/test/CodeGen/NVPTX/address-folder.ll llvm/test/CodeGen/NVPTX/address-folder.mir ################################################################################ diff --git a/llvm/lib/Target/NVPTX/CMakeLists.txt b/llvm/lib/Target/NVPTX/CMakeLists.txt index d2b5083a5f83f..e0e2b1a40bffc 100644 --- a/llvm/lib/Target/NVPTX/CMakeLists.txt +++ b/llvm/lib/Target/NVPTX/CMakeLists.txt @@ -12,7 +12,6 @@ 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 ef447e4ceed6b..5933f74fcb25f 100644 --- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp +++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp @@ -24,7 +24,6 @@ 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 58083e1248bb0..78f4e67455023 100644 --- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.h +++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.h @@ -13,7 +13,6 @@ #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. @@ -23,7 +22,6 @@ // 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 57ed9bb43aeea..08ccc143dfb8a 100644 --- a/llvm/lib/Target/NVPTX/NVPTX.h +++ b/llvm/lib/Target/NVPTX/NVPTX.h @@ -59,7 +59,6 @@ FunctionPass *createNVPTXIRPeepholePass(); MachineFunctionPass *createNVPTXPeephole(); MachineFunctionPass *createNVPTXProxyRegErasurePass(); MachineFunctionPass *createNVPTXForwardParamsPass(); -MachineFunctionPass *createNVPTXAddressFolderPass(); void initializeNVVMReflectLegacyPassPass(PassRegistry &); void initializeGenericToNVVMLegacyPassPass(PassRegistry &); @@ -75,7 +74,6 @@ void initializeNVPTXLowerArgsLegacyPassPass(PassRegistry &); void initializeNVPTXSetByValParamAlignLegacyPassPass(PassRegistry &); void initializeNVPTXProxyRegErasurePass(PassRegistry &); void initializeNVPTXForwardParamsPassPass(PassRegistry &); -void initializeNVPTXAddressFolderPassPass(PassRegistry &); void initializeNVVMIntrRangePass(PassRegistry &); void initializeNVVMReflectPass(PassRegistry &); void initializeNVPTXAAWrapperPassPass(PassRegistry &); @@ -296,7 +294,6 @@ 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 deleted file mode 100644 index 5299dbe686611..0000000000000 --- a/llvm/lib/Target/NVPTX/NVPTXAddressFolder.cpp +++ /dev/null @@ -1,123 +0,0 @@ -//===- 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 361f819a340ae..719be0300940e 100644 --- a/llvm/lib/Target/NVPTX/NVPTXInstrFormats.td +++ b/llvm/lib/Target/NVPTX/NVPTXInstrFormats.td @@ -22,7 +22,6 @@ 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 c9db5c5cec025..5f1c16dc7f0fb 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:$negate), - "@${negate}$a bra \t$target;", + (ins B1:$a, brtarget:$target, BranchFlag:$not), + "@${not}$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 2bc2158fccae1..6054feeedf932 100644 --- a/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp @@ -119,7 +119,6 @@ extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializeNVPTXTarget() { initializeNVPTXLowerAggrCopiesPass(PR); initializeNVPTXProxyRegErasurePass(PR); initializeNVPTXForwardParamsPassPass(PR); - initializeNVPTXAddressFolderPassPass(PR); initializeNVPTXDAGToDAGISelLegacyPass(PR); initializeNVPTXAAWrapperPassPass(PR); initializeNVPTXExternalAAWrapperPass(PR); @@ -422,8 +421,6 @@ 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 deleted file mode 100644 index 239f21594265d..0000000000000 --- a/llvm/test/CodeGen/NVPTX/address-folder.ll +++ /dev/null @@ -1,113 +0,0 @@ -; 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 deleted file mode 100644 index ac50925579506..0000000000000 --- a/llvm/test/CodeGen/NVPTX/address-folder.mir +++ /dev/null @@ -1,60 +0,0 @@ -# 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 -... _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
