[AMDGPU] Support block load/store for CSR #130013
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Add support for using the existing `SCRATCH_STORE_BLOCK` and `SCRATCH_LOAD_BLOCK` instructions for saving and restoring callee-saved VGPRs. This is controlled by a new subtarget feature, `block-vgpr-csr`. It does not include WWM registers - those will be saved and restored individually, just like before. This patch does not change the ABI. Use of this feature may lead to slightly increased stack usage, because the memory is not compacted if certain registers don't have to be transferred (this will happen in practice for calling conventions where the callee and caller saved registers are interleaved in groups of 8). However, if the registers at the end of the block of 32 don't have to be transferred, we don't need to use a whole 128-byte stack slot - we can trim some space off the end of the range. In order to implement this feature, we need to rely less on the target-independent code in the PrologEpilogInserter, so we override several new methods in `SIFrameLowering`. We also add new pseudos, `SI_BLOCK_SPILL_V1024_SAVE/RESTORE`. One peculiarity is that both the SI_BLOCK_V1024_RESTORE pseudo and the SCRATCH_LOAD_BLOCK instructions will have all the registers that are not transferred added as implicit uses. This is done in order to inform LiveRegUnits that those registers are not available before the restore (since we're not really restoring them - so we can't afford to scavenge them). Unfortunately, this trick doesn't work with the save, so before the save all the registers in the block will be unavailable (see the unit test).
|
@llvm/pr-subscribers-backend-amdgpu Author: Diana Picus (rovka) ChangesAdd support for using the existing Use of this feature may lead to slightly increased stack usage, because the memory is not compacted if certain registers don't have to be transferred (this will happen in practice for calling conventions where the callee and caller saved registers are interleaved in groups of 8). However, if the registers at the end of the block of 32 don't have to be transferred, we don't need to use a whole 128-byte stack slot - we can trim some space off the end of the range. In order to implement this feature, we need to rely less on the target-independent code in the PrologEpilogInserter, so we override several new methods in One peculiarity is that both the SI_BLOCK_V1024_RESTORE pseudo and the SCRATCH_LOAD_BLOCK instructions will have all the registers that are not transferred added as implicit uses. This is done in order to inform LiveRegUnits that those registers are not available before the restore (since we're not really restoring them - so we can't afford to scavenge them). Unfortunately, this trick doesn't work with the save, so before the save all the registers in the block will be unavailable (see the unit test). Patch is 78.01 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/130013.diff 18 Files Affected:
diff --git a/llvm/include/llvm/CodeGen/MachineFrameInfo.h b/llvm/include/llvm/CodeGen/MachineFrameInfo.h
index 172c3e8c9a847..e2ee39c42c157 100644
--- a/llvm/include/llvm/CodeGen/MachineFrameInfo.h
+++ b/llvm/include/llvm/CodeGen/MachineFrameInfo.h
@@ -53,6 +53,9 @@ class CalleeSavedInfo {
/// Flag indicating whether the register is spilled to stack or another
/// register.
bool SpilledToReg = false;
+ /// Flag indicating whether this CSI has been handled by the target and can be
+ /// skipped by the generic code in the prolog/epilog inserter.
+ bool IsHandledByTarget = false;
public:
explicit CalleeSavedInfo(MCRegister R, int FI = 0) : Reg(R), FrameIdx(FI) {}
@@ -61,6 +64,7 @@ class CalleeSavedInfo {
MCRegister getReg() const { return Reg; }
int getFrameIdx() const { return FrameIdx; }
MCRegister getDstReg() const { return DstReg; }
+ void setReg(MCRegister R) { Reg = R; }
void setFrameIdx(int FI) {
FrameIdx = FI;
SpilledToReg = false;
@@ -72,6 +76,9 @@ class CalleeSavedInfo {
bool isRestored() const { return Restored; }
void setRestored(bool R) { Restored = R; }
bool isSpilledToReg() const { return SpilledToReg; }
+
+ bool isHandledByTarget() const { return IsHandledByTarget; }
+ void setHandledByTarget() { IsHandledByTarget = true; }
};
/// The MachineFrameInfo class represents an abstract stack frame until
diff --git a/llvm/lib/CodeGen/PrologEpilogInserter.cpp b/llvm/lib/CodeGen/PrologEpilogInserter.cpp
index 9b852c0fd49cf..c74d1b30c24c4 100644
--- a/llvm/lib/CodeGen/PrologEpilogInserter.cpp
+++ b/llvm/lib/CodeGen/PrologEpilogInserter.cpp
@@ -478,7 +478,7 @@ static void assignCalleeSavedSpillSlots(MachineFunction &F,
for (auto &CS : CSI) {
// If the target has spilled this register to another register, we don't
// need to allocate a stack slot.
- if (CS.isSpilledToReg())
+ if (CS.isSpilledToReg() || CS.isHandledByTarget())
continue;
MCRegister Reg = CS.getReg();
@@ -604,6 +604,8 @@ static void insertCSRSaves(MachineBasicBlock &SaveBlock,
MachineBasicBlock::iterator I = SaveBlock.begin();
if (!TFI->spillCalleeSavedRegisters(SaveBlock, I, CSI, TRI)) {
for (const CalleeSavedInfo &CS : CSI) {
+ if (CS.isHandledByTarget())
+ continue;
// Insert the spill to the stack frame.
MCRegister Reg = CS.getReg();
@@ -634,6 +636,9 @@ static void insertCSRRestores(MachineBasicBlock &RestoreBlock,
if (!TFI->restoreCalleeSavedRegisters(RestoreBlock, I, CSI, TRI)) {
for (const CalleeSavedInfo &CI : reverse(CSI)) {
+ if (CI.isHandledByTarget())
+ continue;
+
MCRegister Reg = CI.getReg();
if (CI.isSpilledToReg()) {
BuildMI(RestoreBlock, I, DebugLoc(), TII.get(TargetOpcode::COPY), Reg)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index effc8d2ed6b49..4c9268d6ff6a6 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.td
@@ -1239,6 +1239,14 @@ def FeatureXF32Insts : SubtargetFeature<"xf32-insts",
"v_mfma_f32_16x16x8_xf32 and v_mfma_f32_32x32x4_xf32"
>;
+// Enable the use of SCRATCH_STORE/LOAD_BLOCK instructions for saving and
+// restoring the callee-saved registers.
+def FeatureUseBlockVGPROpsForCSR : SubtargetFeature<"block-vgpr-csr",
+ "UseBlockVGPROpsForCSR",
+ "true",
+ "Use block load/store for VGPR callee saved registers"
+>;
+
// Dummy feature used to disable assembler instructions.
def FeatureDisable : SubtargetFeature<"",
"FeatureDisable","true",
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp b/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp
index 895d1e77bf1c4..60b8c5a7a2251 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp
@@ -18,6 +18,7 @@
#include "AMDGPUMachineFunction.h"
#include "MCTargetDesc/AMDGPUInstPrinter.h"
#include "MCTargetDesc/AMDGPUMCTargetDesc.h"
+#include "SIMachineFunctionInfo.h"
#include "llvm/CodeGen/MachineBasicBlock.h"
#include "llvm/CodeGen/MachineInstr.h"
#include "llvm/IR/Constants.h"
@@ -239,6 +240,36 @@ const MCExpr *AMDGPUAsmPrinter::lowerConstant(const Constant *CV) {
return AsmPrinter::lowerConstant(CV);
}
+static void emitVGPRBlockComment(const MachineInstr *MI, MCStreamer &OS) {
+ // The instruction will only transfer a subset of the registers in the block,
+ // based on the mask that is stored in m0. We could search for the instruction
+ // that sets m0, but most of the time we'll already have the mask stored in
+ // the machine function info. Try to use that. This assumes that we only use
+ // block loads/stores for CSR spills.
+ const MachineFunction *MF = MI->getParent()->getParent();
+ const SIMachineFunctionInfo *MFI = MF->getInfo<SIMachineFunctionInfo>();
+ const TargetRegisterInfo &TRI = *MF->getSubtarget().getRegisterInfo();
+ const SIInstrInfo *TII = MF->getSubtarget<GCNSubtarget>().getInstrInfo();
+
+ Register RegBlock =
+ TII->getNamedOperand(*MI, MI->mayLoad() ? AMDGPU::OpName::vdst
+ : AMDGPU::OpName::vdata)
+ ->getReg();
+ Register FirstRegInBlock = TRI.getSubReg(RegBlock, AMDGPU::sub0);
+ uint32_t Mask = MFI->getMaskForVGPRBlockOps(RegBlock);
+
+ SmallString<512> TransferredRegs;
+ for (unsigned I = 0; I < 32; ++I) {
+ if (Mask & (1 << I)) {
+ (llvm::Twine(" ") + TRI.getName(FirstRegInBlock + I))
+ .toVector(TransferredRegs);
+ }
+ }
+
+ if (!TransferredRegs.empty())
+ OS.emitRawComment(" transferring at most " + TransferredRegs);
+}
+
void AMDGPUAsmPrinter::emitInstruction(const MachineInstr *MI) {
// FIXME: Enable feature predicate checks once all the test pass.
// AMDGPU_MC::verifyInstructionPredicates(MI->getOpcode(),
@@ -327,6 +358,10 @@ void AMDGPUAsmPrinter::emitInstruction(const MachineInstr *MI) {
return;
}
+ if (STI.getInstrInfo()->isBlockLoadStore(MI->getOpcode()))
+ if (isVerbose())
+ emitVGPRBlockComment(MI, *OutStreamer);
+
MCInst TmpInst;
MCInstLowering.lower(MI, TmpInst);
EmitToStreamer(*OutStreamer, TmpInst);
diff --git a/llvm/lib/Target/AMDGPU/GCNSubtarget.h b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
index 6664a70572ded..033f2870fc6a0 100644
--- a/llvm/lib/Target/AMDGPU/GCNSubtarget.h
+++ b/llvm/lib/Target/AMDGPU/GCNSubtarget.h
@@ -254,6 +254,7 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
bool HasMinimum3Maximum3PKF16 = false;
bool RequiresCOV6 = false;
+ bool UseBlockVGPROpsForCSR = false;
// Dummy feature to use for assembler in tablegen.
bool FeatureDisable = false;
@@ -1265,6 +1266,8 @@ class GCNSubtarget final : public AMDGPUGenSubtargetInfo,
bool requiresCodeObjectV6() const { return RequiresCOV6; }
+ bool useVGPRBlockOpsForCSR() const { return UseBlockVGPROpsForCSR; }
+
bool hasVALUMaskWriteHazard() const { return getGeneration() == GFX11; }
bool hasVALUReadSGPRHazard() const { return getGeneration() == GFX12; }
diff --git a/llvm/lib/Target/AMDGPU/SIFrameLowering.cpp b/llvm/lib/Target/AMDGPU/SIFrameLowering.cpp
index 97736e2410c18..6c9cfba3196bf 100644
--- a/llvm/lib/Target/AMDGPU/SIFrameLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIFrameLowering.cpp
@@ -1694,6 +1694,110 @@ void SIFrameLowering::determineCalleeSavesSGPR(MachineFunction &MF,
}
}
+static void assignSlotsUsingVGPRBlocks(MachineFunction &MF,
+ const GCNSubtarget &ST,
+ const TargetRegisterInfo *TRI,
+ std::vector<CalleeSavedInfo> &CSI,
+ unsigned &MinCSFrameIndex,
+ unsigned &MaxCSFrameIndex) {
+ SIMachineFunctionInfo *FuncInfo = MF.getInfo<SIMachineFunctionInfo>();
+ MachineFrameInfo &MFI = MF.getFrameInfo();
+ const SIInstrInfo *TII = ST.getInstrInfo();
+ const SIRegisterInfo *MRI = ST.getRegisterInfo();
+
+ assert(std::is_sorted(CSI.begin(), CSI.end(),
+ [](const CalleeSavedInfo &A, const CalleeSavedInfo &B) {
+ return A.getReg() < B.getReg();
+ }) &&
+ "Callee saved registers not sorted");
+
+ auto CanUseBlockOps = [&](const CalleeSavedInfo &CSI) {
+ return !CSI.isSpilledToReg() &&
+ MRI->isVGPR(MF.getRegInfo(), CSI.getReg()) &&
+ !FuncInfo->isWWMReservedRegister(CSI.getReg());
+ };
+
+ auto CSEnd = CSI.end();
+ for (auto CSIt = CSI.begin(); CSIt != CSEnd; ++CSIt) {
+ Register Reg = CSIt->getReg();
+ if (!CanUseBlockOps(*CSIt))
+ continue;
+
+ // Find all the regs that will fit in a 32-bit block starting at the current
+ // reg and build the mask. It should have 1 for every register that's
+ // included, with the current register as the least significant bit.
+ uint32_t Mask = 1;
+ CSEnd = std::remove_if(
+ CSIt + 1, CSEnd, [&](const CalleeSavedInfo &CSI) -> bool {
+ if (CanUseBlockOps(CSI) && CSI.getReg() < Reg + 32) {
+ Mask |= 1 << (CSI.getReg() - Reg);
+ return true;
+ } else {
+ return false;
+ }
+ });
+
+ const TargetRegisterClass *BlockRegClass =
+ TII->getRegClassForBlockOp(TRI, MF);
+ Register RegBlock =
+ MRI->getMatchingSuperReg(Reg, AMDGPU::sub0, BlockRegClass);
+ if (!RegBlock) {
+ // We couldn't find a super register for the block. This can happen if
+ // the register we started with is too high (e.g. v232 if the maximum is
+ // v255). We therefore try to get the last register block and figure out
+ // the mask from there.
+ Register LastBlockStart =
+ AMDGPU::VGPR0 + alignDown(Reg - AMDGPU::VGPR0, 32);
+ RegBlock =
+ MRI->getMatchingSuperReg(LastBlockStart, AMDGPU::sub0, BlockRegClass);
+ assert(RegBlock && MRI->isSubRegister(RegBlock, Reg) &&
+ "Couldn't find super register");
+ int RegDelta = Reg - LastBlockStart;
+ assert(RegDelta > 0 && llvm::countl_zero(Mask) >= RegDelta &&
+ "Bad shift amount");
+ Mask <<= RegDelta;
+ }
+
+ FuncInfo->setMaskForVGPRBlockOps(RegBlock, Mask);
+
+ // The stack objects can be a bit smaller than the register block if we know
+ // some of the high bits of Mask are 0. This may happen often with calling
+ // conventions where the caller and callee-saved VGPRs are interleaved at
+ // a small boundary (e.g. 8 or 16).
+ int UnusedBits = llvm::countl_zero(Mask);
+ unsigned BlockSize = MRI->getSpillSize(*BlockRegClass) - UnusedBits * 4;
+ int FrameIdx =
+ MFI.CreateStackObject(BlockSize, MRI->getSpillAlign(*BlockRegClass),
+ /*isSpillSlot=*/true);
+ if ((unsigned)FrameIdx < MinCSFrameIndex)
+ MinCSFrameIndex = FrameIdx;
+ if ((unsigned)FrameIdx > MaxCSFrameIndex)
+ MaxCSFrameIndex = FrameIdx;
+
+ CSIt->setFrameIdx(FrameIdx);
+ CSIt->setReg(RegBlock);
+ CSIt->setHandledByTarget();
+ }
+ CSI.erase(CSEnd, CSI.end());
+}
+
+bool SIFrameLowering::assignCalleeSavedSpillSlots(
+ MachineFunction &MF, const TargetRegisterInfo *TRI,
+ std::vector<CalleeSavedInfo> &CSI, unsigned &MinCSFrameIndex,
+ unsigned &MaxCSFrameIndex) const {
+ if (CSI.empty())
+ return true; // Early exit if no callee saved registers are modified!
+
+ const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
+ bool UseVGPRBlocks = ST.useVGPRBlockOpsForCSR();
+
+ if (UseVGPRBlocks)
+ assignSlotsUsingVGPRBlocks(MF, ST, TRI, CSI, MinCSFrameIndex,
+ MaxCSFrameIndex);
+
+ return assignCalleeSavedSpillSlots(MF, TRI, CSI);
+}
+
bool SIFrameLowering::assignCalleeSavedSpillSlots(
MachineFunction &MF, const TargetRegisterInfo *TRI,
std::vector<CalleeSavedInfo> &CSI) const {
@@ -1763,6 +1867,101 @@ bool SIFrameLowering::allocateScavengingFrameIndexesNearIncomingSP(
return true;
}
+bool SIFrameLowering::spillCalleeSavedRegisters(
+ MachineBasicBlock &MBB, MachineBasicBlock::iterator MI,
+ ArrayRef<CalleeSavedInfo> CSI, const TargetRegisterInfo *TRI) const {
+ MachineFunction *MF = MBB.getParent();
+ const GCNSubtarget &ST = MF->getSubtarget<GCNSubtarget>();
+ if (!ST.useVGPRBlockOpsForCSR())
+ return false;
+
+ MachineFrameInfo &FrameInfo = MF->getFrameInfo();
+ SIMachineFunctionInfo *MFI = MF->getInfo<SIMachineFunctionInfo>();
+ const SIInstrInfo *TII = ST.getInstrInfo();
+ SIMachineFunctionInfo *FuncInfo = MF->getInfo<SIMachineFunctionInfo>();
+
+ for (const CalleeSavedInfo &CS : CSI) {
+ Register Reg = CS.getReg();
+ if (!CS.isHandledByTarget())
+ continue;
+
+ // Build a scratch block store.
+ uint32_t Mask = FuncInfo->getMaskForVGPRBlockOps(Reg);
+ int FrameIndex = CS.getFrameIdx();
+ MachinePointerInfo PtrInfo =
+ MachinePointerInfo::getFixedStack(*MF, FrameIndex);
+ MachineMemOperand *MMO =
+ MF->getMachineMemOperand(PtrInfo, MachineMemOperand::MOStore,
+ FrameInfo.getObjectSize(FrameIndex),
+ FrameInfo.getObjectAlign(FrameIndex));
+
+ BuildMI(MBB, MI, MI->getDebugLoc(),
+ TII->get(AMDGPU::SI_BLOCK_SPILL_V1024_SAVE))
+ .addReg(Reg, getKillRegState(false))
+ .addFrameIndex(FrameIndex)
+ .addReg(MFI->getStackPtrOffsetReg())
+ .addImm(0)
+ .addImm(Mask)
+ .addMemOperand(MMO);
+
+ FuncInfo->setHasSpilledVGPRs();
+
+ // Add the register to the liveins. This is necessary because if any of the
+ // VGPRs in the register block is reserved (e.g. if it's a WWM register),
+ // then the whole block will be marked as reserved and `updateLiveness` will
+ // skip it.
+ MBB.addLiveIn(Reg);
+ }
+
+ return false;
+}
+
+bool SIFrameLowering::restoreCalleeSavedRegisters(
+ MachineBasicBlock &MBB, MachineBasicBlock::iterator MI,
+ MutableArrayRef<CalleeSavedInfo> CSI, const TargetRegisterInfo *TRI) const {
+ MachineFunction *MF = MBB.getParent();
+ const GCNSubtarget &ST = MF->getSubtarget<GCNSubtarget>();
+ if (!ST.useVGPRBlockOpsForCSR())
+ return false;
+
+ SIMachineFunctionInfo *FuncInfo = MF->getInfo<SIMachineFunctionInfo>();
+ MachineFrameInfo &MFI = MF->getFrameInfo();
+ const SIInstrInfo *TII = ST.getInstrInfo();
+ const SIRegisterInfo *SITRI = static_cast<const SIRegisterInfo *>(TRI);
+ for (const CalleeSavedInfo &CS : reverse(CSI)) {
+ if (!CS.isHandledByTarget())
+ continue;
+
+ // Build a scratch block load.
+ Register Reg = CS.getReg();
+ uint32_t Mask = FuncInfo->getMaskForVGPRBlockOps(Reg);
+ int FrameIndex = CS.getFrameIdx();
+ MachinePointerInfo PtrInfo =
+ MachinePointerInfo::getFixedStack(*MF, FrameIndex);
+ MachineMemOperand *MMO = MF->getMachineMemOperand(
+ PtrInfo, MachineMemOperand::MOLoad, MFI.getObjectSize(FrameIndex),
+ MFI.getObjectAlign(FrameIndex));
+
+ auto MIB = BuildMI(MBB, MI, MI->getDebugLoc(),
+ TII->get(AMDGPU::SI_BLOCK_SPILL_V1024_RESTORE), Reg)
+ .addFrameIndex(FrameIndex)
+ .addReg(FuncInfo->getStackPtrOffsetReg())
+ .addImm(0)
+ .addImm(Mask)
+ .addMemOperand(MMO);
+ SITRI->addImplicitUsesForBlockCSRLoad(MIB, Reg);
+
+ // Add the register to the liveins. This is necessary because if any of the
+ // VGPRs in the register block is reserved (e.g. if it's a WWM register),
+ // then the whole block will be marked as reserved and `updateLiveness` will
+ // skip it.
+ if (!MBB.isLiveIn(Reg))
+ MBB.addLiveIn(Reg);
+ }
+
+ return false;
+}
+
MachineBasicBlock::iterator SIFrameLowering::eliminateCallFramePseudoInstr(
MachineFunction &MF,
MachineBasicBlock &MBB,
diff --git a/llvm/lib/Target/AMDGPU/SIFrameLowering.h b/llvm/lib/Target/AMDGPU/SIFrameLowering.h
index 938c75099a3bc..e20ded586214e 100644
--- a/llvm/lib/Target/AMDGPU/SIFrameLowering.h
+++ b/llvm/lib/Target/AMDGPU/SIFrameLowering.h
@@ -49,6 +49,23 @@ class SIFrameLowering final : public AMDGPUFrameLowering {
const TargetRegisterInfo *TRI,
std::vector<CalleeSavedInfo> &CSI) const override;
+ bool assignCalleeSavedSpillSlots(MachineFunction &MF,
+ const TargetRegisterInfo *TRI,
+ std::vector<CalleeSavedInfo> &CSI,
+ unsigned &MinCSFrameIndex,
+ unsigned &MaxCSFrameIndex) const override;
+
+ bool spillCalleeSavedRegisters(MachineBasicBlock &MBB,
+ MachineBasicBlock::iterator MI,
+ ArrayRef<CalleeSavedInfo> CSI,
+ const TargetRegisterInfo *TRI) const override;
+
+ bool
+ restoreCalleeSavedRegisters(MachineBasicBlock &MBB,
+ MachineBasicBlock::iterator MI,
+ MutableArrayRef<CalleeSavedInfo> CSI,
+ const TargetRegisterInfo *TRI) const override;
+
bool allocateScavengingFrameIndexesNearIncomingSP(
const MachineFunction &MF) const override;
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
index e1fb4c9db0a84..535ebe3d30739 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
@@ -5831,6 +5831,16 @@ const TargetRegisterClass *SIInstrInfo::getRegClass(const MCInstrDesc &TID,
IsAllocatable);
}
+const TargetRegisterClass *
+SIInstrInfo::getRegClassForBlockOp(const TargetRegisterInfo *TRI,
+ const MachineFunction &MF) const {
+ const MCInstrDesc &ScratchStoreBlockOp =
+ get(AMDGPU::SCRATCH_STORE_BLOCK_SADDR);
+ int VDataIdx = AMDGPU::getNamedOperandIdx(ScratchStoreBlockOp.getOpcode(),
+ AMDGPU::OpName::vdata);
+ return getRegClass(ScratchStoreBlockOp, VDataIdx, TRI, MF);
+}
+
const TargetRegisterClass *SIInstrInfo::getOpRegClass(const MachineInstr &MI,
unsigned OpNo) const {
const MachineRegisterInfo &MRI = MI.getParent()->getParent()->getRegInfo();
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.h b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
index 79ef1432d512a..82648f14b252c 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
@@ -663,6 +663,18 @@ class SIInstrInfo final : public AMDGPUGenInstrInfo {
return get(Opcode).TSFlags & SIInstrFlags::FLAT;
}
+ static bool isBlockLoadStore(uint16_t Opcode) {
+ switch (Opcode) {
+ case AMDGPU::SI_BLOCK_SPILL_V1024_SAVE:
+ case AMDGPU::SI_BLOCK_SPILL_V1024_RESTORE:
+ case AMDGPU::SCRATCH_STORE_BLOCK_SADDR:
+ case AMDGPU::SCRATCH_LOAD_BLOCK_SADDR:
+ return true;
+ default:
+ return false;
+ }
+ }
+
static bool isEXP(const MachineInstr &MI) {
return MI.getDesc().TSFlags & SIInstrFlags::EXP;
}
@@ -1448,6 +1460,10 @@ class SIInstrInfo final : public AMDGPUGenInstrInfo {
const MachineFunction &MF)
const override;
+ const TargetRegisterClass *
+ getRegClassForBlockOp(const TargetRegisterInfo *TRI,
+ const MachineFunction &MF) const;
+
void fixImplicitOperands(MachineInstr &MI) const;
MachineInstr *foldMemoryOperandImpl(MachineFunction &MF, MachineInstr &MI,
diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td
index 63f66023837a2..f9ad2d5621120 100644
--- a/llvm/lib/Target/AMDGPU/SIInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SIInstructions.td
@@ -973,13 +973,16 @@ def SI_RESTORE_S32_FROM_VGPR : PseudoInstSI <(outs SReg_32:$sdst),
// VGPR or AGPR spill instructions. In case of AGPR spilling a temp register
// needs to be used and an extra instruction to move between VGPR and AGPR.
// UsesTmp adds to the total size of an expanded spill in this case.
-multiclass SI_SPILL_VGPR <RegisterClass vgpr_class, bit UsesTmp = 0> {
+multiclass SI_SPILL_VGPR <RegisterClass vgpr_class,
+ ...
[truncated]
|
|
Ping? |
perlfu
reviewed
Mar 10, 2025
| CSEnd = std::remove_if( | ||
| CSIt + 1, CSEnd, [&](const CalleeSavedInfo &CSI) -> bool { | ||
| if (CanUseBlockOps(CSI) && CSI.getReg() < Reg + 32) { | ||
| Mask |= 1 << (CSI.getReg() - Reg); |
There was a problem hiding this comment.
Seems a little odd to build the Mask through the predicate, but I guess it will be deterministic even if for example the predicate function was called multiple times on the same element.
slinder1
reviewed
Mar 11, 2025
arsenm
reviewed
Mar 12, 2025
| Register FirstRegInBlock = TRI.getSubReg(RegBlock, AMDGPU::sub0); | ||
| uint32_t Mask = MFI->getMaskForVGPRBlockOps(RegBlock); | ||
|
|
||
| SmallString<512> TransferredRegs; |
There was a problem hiding this comment.
You can avoid the temporary SmallString, emitRawComment uses Twine anyway
There was a problem hiding this comment.
Wasn't Twine only meant to be used in function arguments?
Or did you mean to emit one comment for each transferred register?
qiaojbao
added a commit
to GPUOpen-Drivers/llvm-project
that referenced
this pull request
Mar 13, 2025
qiaojbao
added a commit
to GPUOpen-Drivers/llvm-project
that referenced
this pull request
Mar 13, 2025
|
Ping. I refactored the PEI mechanisms a bit. |
arsenm
reviewed
Mar 18, 2025
Comment on lines
582
to
584
| // This is only useful during prolog/epilog insertion, so it doesn't need to | ||
| // be serialized. | ||
| IndexedMap<uint32_t, VGPRBlock2IndexFunctor> MaskForVGPRBlockOps; |
There was a problem hiding this comment.
I don't follow why this needs to be here. If it doesn't need serialization, it is pass local state that should not be here. This is only used on instructions directly inserted in PEI. Why isn't this directly represented with a set of implicit uses (or a regmask) inline on the instruction? This seems like an avoidable intermediate state
There was a problem hiding this comment.
Actually that comment is obsolete, we also use this when printing out the assembly (so we can put the comment about which regs might be getting spilled). It's difficult to add this to the instructions because the mask is first computed in assignCalleeSavedSpillSlots, which doesn't create any instructions yet.
|
Ping |
cdevadas
approved these changes
Apr 8, 2025
|
There haven't been any new concerns for the past 2 weeks, so I'm going to merge this. |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/16/builds/17800 Here is the relevant piece of the build log for the reference |
RKSimon
reviewed
Apr 23, 2025
| @@ -388,6 +388,16 @@ class PrologEpilogSGPRSaveRestoreInfo { | |||
| SGPRSaveKind getKind() const { return Kind; } | |||
| }; | |||
|
|
|||
| const MCRegister FirstVGPRBlock = AMDGPU::VReg_1024RegClass.getRegister(0); | |||
There was a problem hiding this comment.
I'm seeing initialization ordering failures from this - can't we just put it inside the VGPRBlock2IndexFunctor operator?
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/186/builds/8439 Here is the relevant piece of the build log for the reference |
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/187/builds/5564 Here is the relevant piece of the build log for the reference |
rovka
added a commit
that referenced
this pull request
Apr 23, 2025
Reverts #130013 due to failures with expensive checks on.
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/146/builds/2769 Here is the relevant piece of the build log for the reference |
rovka
added a commit
that referenced
this pull request
Apr 25, 2025
Add support for using the existing SCRATCH_STORE_BLOCK and SCRATCH_LOAD_BLOCK instructions for saving and restoring callee-saved VGPRs. This is controlled by a new subtarget feature, block-vgpr-csr. It does not include WWM registers - those will be saved and restored individually, just like before. This patch does not change the ABI. Use of this feature may lead to slightly increased stack usage, because the memory is not compacted if certain registers don't have to be transferred (this will happen in practice for calling conventions where the callee and caller saved registers are interleaved in groups of 8). However, if the registers at the end of the block of 32 don't have to be transferred, we don't need to use a whole 128-byte stack slot - we can trim some space off the end of the range. In order to implement this feature, we need to rely less on the target-independent code in the PrologEpilogInserter, so we override several new methods in SIFrameLowering. We also add new pseudos, SI_BLOCK_SPILL_V1024_SAVE/RESTORE. One peculiarity is that both the SI_BLOCK_V1024_RESTORE pseudo and the SCRATCH_LOAD_BLOCK instructions will have all the registers that are not transferred added as implicit uses. This is done in order to inform LiveRegUnits that those registers are not available before the restore (since we're not really restoring them - so we can't afford to scavenge them). Unfortunately, this trick doesn't work with the save, so before the save all the registers in the block will be unavailable (see the unit test). This was reverted due to failures in the builds with expensive checks on, now fixed by always updating LiveIntervals and SlotIndexes in SILowerSGPRSpills.
searlmc1
pushed a commit
to ROCm/llvm-project
that referenced
this pull request
Apr 25, 2025
…lvm#137169)" needs to be integrated into the downstream divergence. This reverts commit 5bad5d8.
searlmc1
pushed a commit
to ROCm/llvm-project
that referenced
this pull request
Apr 29, 2025
…llvm#137169)" This reverts commit c076481.
llvm-sync bot
pushed a commit
to arm/arm-toolchain
that referenced
this pull request
May 6, 2025
Reverts llvm/llvm-project#130013 due to failures with expensive checks on.
IanWood1
pushed a commit
to IanWood1/llvm-project
that referenced
this pull request
May 6, 2025
Add support for using the existing `SCRATCH_STORE_BLOCK` and `SCRATCH_LOAD_BLOCK` instructions for saving and restoring callee-saved VGPRs. This is controlled by a new subtarget feature, `block-vgpr-csr`. It does not include WWM registers - those will be saved and restored individually, just like before. This patch does not change the ABI. Use of this feature may lead to slightly increased stack usage, because the memory is not compacted if certain registers don't have to be transferred (this will happen in practice for calling conventions where the callee and caller saved registers are interleaved in groups of 8). However, if the registers at the end of the block of 32 don't have to be transferred, we don't need to use a whole 128-byte stack slot - we can trim some space off the end of the range. In order to implement this feature, we need to rely less on the target-independent code in the PrologEpilogInserter, so we override several new methods in `SIFrameLowering`. We also add new pseudos, `SI_BLOCK_SPILL_V1024_SAVE/RESTORE`. One peculiarity is that both the SI_BLOCK_V1024_RESTORE pseudo and the SCRATCH_LOAD_BLOCK instructions will have all the registers that are not transferred added as implicit uses. This is done in order to inform LiveRegUnits that those registers are not available before the restore (since we're not really restoring them - so we can't afford to scavenge them). Unfortunately, this trick doesn't work with the save, so before the save all the registers in the block will be unavailable (see the unit test).
IanWood1
pushed a commit
to IanWood1/llvm-project
that referenced
this pull request
May 6, 2025
Reverts llvm#130013 due to failures with expensive checks on.
IanWood1
pushed a commit
to IanWood1/llvm-project
that referenced
this pull request
May 6, 2025
…7169) Add support for using the existing SCRATCH_STORE_BLOCK and SCRATCH_LOAD_BLOCK instructions for saving and restoring callee-saved VGPRs. This is controlled by a new subtarget feature, block-vgpr-csr. It does not include WWM registers - those will be saved and restored individually, just like before. This patch does not change the ABI. Use of this feature may lead to slightly increased stack usage, because the memory is not compacted if certain registers don't have to be transferred (this will happen in practice for calling conventions where the callee and caller saved registers are interleaved in groups of 8). However, if the registers at the end of the block of 32 don't have to be transferred, we don't need to use a whole 128-byte stack slot - we can trim some space off the end of the range. In order to implement this feature, we need to rely less on the target-independent code in the PrologEpilogInserter, so we override several new methods in SIFrameLowering. We also add new pseudos, SI_BLOCK_SPILL_V1024_SAVE/RESTORE. One peculiarity is that both the SI_BLOCK_V1024_RESTORE pseudo and the SCRATCH_LOAD_BLOCK instructions will have all the registers that are not transferred added as implicit uses. This is done in order to inform LiveRegUnits that those registers are not available before the restore (since we're not really restoring them - so we can't afford to scavenge them). Unfortunately, this trick doesn't work with the save, so before the save all the registers in the block will be unavailable (see the unit test). This was reverted due to failures in the builds with expensive checks on, now fixed by always updating LiveIntervals and SlotIndexes in SILowerSGPRSpills.
IanWood1
pushed a commit
to IanWood1/llvm-project
that referenced
this pull request
May 6, 2025
Add support for using the existing `SCRATCH_STORE_BLOCK` and `SCRATCH_LOAD_BLOCK` instructions for saving and restoring callee-saved VGPRs. This is controlled by a new subtarget feature, `block-vgpr-csr`. It does not include WWM registers - those will be saved and restored individually, just like before. This patch does not change the ABI. Use of this feature may lead to slightly increased stack usage, because the memory is not compacted if certain registers don't have to be transferred (this will happen in practice for calling conventions where the callee and caller saved registers are interleaved in groups of 8). However, if the registers at the end of the block of 32 don't have to be transferred, we don't need to use a whole 128-byte stack slot - we can trim some space off the end of the range. In order to implement this feature, we need to rely less on the target-independent code in the PrologEpilogInserter, so we override several new methods in `SIFrameLowering`. We also add new pseudos, `SI_BLOCK_SPILL_V1024_SAVE/RESTORE`. One peculiarity is that both the SI_BLOCK_V1024_RESTORE pseudo and the SCRATCH_LOAD_BLOCK instructions will have all the registers that are not transferred added as implicit uses. This is done in order to inform LiveRegUnits that those registers are not available before the restore (since we're not really restoring them - so we can't afford to scavenge them). Unfortunately, this trick doesn't work with the save, so before the save all the registers in the block will be unavailable (see the unit test).
IanWood1
pushed a commit
to IanWood1/llvm-project
that referenced
this pull request
May 6, 2025
Reverts llvm#130013 due to failures with expensive checks on.
IanWood1
pushed a commit
to IanWood1/llvm-project
that referenced
this pull request
May 6, 2025
…7169) Add support for using the existing SCRATCH_STORE_BLOCK and SCRATCH_LOAD_BLOCK instructions for saving and restoring callee-saved VGPRs. This is controlled by a new subtarget feature, block-vgpr-csr. It does not include WWM registers - those will be saved and restored individually, just like before. This patch does not change the ABI. Use of this feature may lead to slightly increased stack usage, because the memory is not compacted if certain registers don't have to be transferred (this will happen in practice for calling conventions where the callee and caller saved registers are interleaved in groups of 8). However, if the registers at the end of the block of 32 don't have to be transferred, we don't need to use a whole 128-byte stack slot - we can trim some space off the end of the range. In order to implement this feature, we need to rely less on the target-independent code in the PrologEpilogInserter, so we override several new methods in SIFrameLowering. We also add new pseudos, SI_BLOCK_SPILL_V1024_SAVE/RESTORE. One peculiarity is that both the SI_BLOCK_V1024_RESTORE pseudo and the SCRATCH_LOAD_BLOCK instructions will have all the registers that are not transferred added as implicit uses. This is done in order to inform LiveRegUnits that those registers are not available before the restore (since we're not really restoring them - so we can't afford to scavenge them). Unfortunately, this trick doesn't work with the save, so before the save all the registers in the block will be unavailable (see the unit test). This was reverted due to failures in the builds with expensive checks on, now fixed by always updating LiveIntervals and SlotIndexes in SILowerSGPRSpills.
IanWood1
pushed a commit
to IanWood1/llvm-project
that referenced
this pull request
May 6, 2025
Add support for using the existing `SCRATCH_STORE_BLOCK` and `SCRATCH_LOAD_BLOCK` instructions for saving and restoring callee-saved VGPRs. This is controlled by a new subtarget feature, `block-vgpr-csr`. It does not include WWM registers - those will be saved and restored individually, just like before. This patch does not change the ABI. Use of this feature may lead to slightly increased stack usage, because the memory is not compacted if certain registers don't have to be transferred (this will happen in practice for calling conventions where the callee and caller saved registers are interleaved in groups of 8). However, if the registers at the end of the block of 32 don't have to be transferred, we don't need to use a whole 128-byte stack slot - we can trim some space off the end of the range. In order to implement this feature, we need to rely less on the target-independent code in the PrologEpilogInserter, so we override several new methods in `SIFrameLowering`. We also add new pseudos, `SI_BLOCK_SPILL_V1024_SAVE/RESTORE`. One peculiarity is that both the SI_BLOCK_V1024_RESTORE pseudo and the SCRATCH_LOAD_BLOCK instructions will have all the registers that are not transferred added as implicit uses. This is done in order to inform LiveRegUnits that those registers are not available before the restore (since we're not really restoring them - so we can't afford to scavenge them). Unfortunately, this trick doesn't work with the save, so before the save all the registers in the block will be unavailable (see the unit test).
IanWood1
pushed a commit
to IanWood1/llvm-project
that referenced
this pull request
May 6, 2025
Reverts llvm#130013 due to failures with expensive checks on.
IanWood1
pushed a commit
to IanWood1/llvm-project
that referenced
this pull request
May 6, 2025
…7169) Add support for using the existing SCRATCH_STORE_BLOCK and SCRATCH_LOAD_BLOCK instructions for saving and restoring callee-saved VGPRs. This is controlled by a new subtarget feature, block-vgpr-csr. It does not include WWM registers - those will be saved and restored individually, just like before. This patch does not change the ABI. Use of this feature may lead to slightly increased stack usage, because the memory is not compacted if certain registers don't have to be transferred (this will happen in practice for calling conventions where the callee and caller saved registers are interleaved in groups of 8). However, if the registers at the end of the block of 32 don't have to be transferred, we don't need to use a whole 128-byte stack slot - we can trim some space off the end of the range. In order to implement this feature, we need to rely less on the target-independent code in the PrologEpilogInserter, so we override several new methods in SIFrameLowering. We also add new pseudos, SI_BLOCK_SPILL_V1024_SAVE/RESTORE. One peculiarity is that both the SI_BLOCK_V1024_RESTORE pseudo and the SCRATCH_LOAD_BLOCK instructions will have all the registers that are not transferred added as implicit uses. This is done in order to inform LiveRegUnits that those registers are not available before the restore (since we're not really restoring them - so we can't afford to scavenge them). Unfortunately, this trick doesn't work with the save, so before the save all the registers in the block will be unavailable (see the unit test). This was reverted due to failures in the builds with expensive checks on, now fixed by always updating LiveIntervals and SlotIndexes in SILowerSGPRSpills.
Ankur-0429
pushed a commit
to Ankur-0429/llvm-project
that referenced
this pull request
May 9, 2025
…7169) Add support for using the existing SCRATCH_STORE_BLOCK and SCRATCH_LOAD_BLOCK instructions for saving and restoring callee-saved VGPRs. This is controlled by a new subtarget feature, block-vgpr-csr. It does not include WWM registers - those will be saved and restored individually, just like before. This patch does not change the ABI. Use of this feature may lead to slightly increased stack usage, because the memory is not compacted if certain registers don't have to be transferred (this will happen in practice for calling conventions where the callee and caller saved registers are interleaved in groups of 8). However, if the registers at the end of the block of 32 don't have to be transferred, we don't need to use a whole 128-byte stack slot - we can trim some space off the end of the range. In order to implement this feature, we need to rely less on the target-independent code in the PrologEpilogInserter, so we override several new methods in SIFrameLowering. We also add new pseudos, SI_BLOCK_SPILL_V1024_SAVE/RESTORE. One peculiarity is that both the SI_BLOCK_V1024_RESTORE pseudo and the SCRATCH_LOAD_BLOCK instructions will have all the registers that are not transferred added as implicit uses. This is done in order to inform LiveRegUnits that those registers are not available before the restore (since we're not really restoring them - so we can't afford to scavenge them). Unfortunately, this trick doesn't work with the save, so before the save all the registers in the block will be unavailable (see the unit test). This was reverted due to failures in the builds with expensive checks on, now fixed by always updating LiveIntervals and SlotIndexes in SILowerSGPRSpills.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Add support for using the existing
SCRATCH_STORE_BLOCKandSCRATCH_LOAD_BLOCKinstructions for saving and restoring callee-saved VGPRs. This is controlled by a new subtarget feature,block-vgpr-csr. It does not include WWM registers - those will be saved and restored individually, just like before. This patch does not change the ABI.Use of this feature may lead to slightly increased stack usage, because the memory is not compacted if certain registers don't have to be transferred (this will happen in practice for calling conventions where the callee and caller saved registers are interleaved in groups of 8). However, if the registers at the end of the block of 32 don't have to be transferred, we don't need to use a whole 128-byte stack slot - we can trim some space off the end of the range.
In order to implement this feature, we need to rely less on the target-independent code in the PrologEpilogInserter, so we override several new methods in
SIFrameLowering. We also add new pseudos,SI_BLOCK_SPILL_V1024_SAVE/RESTORE.One peculiarity is that both the SI_BLOCK_V1024_RESTORE pseudo and the SCRATCH_LOAD_BLOCK instructions will have all the registers that are not transferred added as implicit uses. This is done in order to inform LiveRegUnits that those registers are not available before the restore (since we're not really restoring them - so we can't afford to scavenge them). Unfortunately, this trick doesn't work with the save, so before the save all the registers in the block will be unavailable (see the unit test).