[llvm] [AMDGPU] Post-RA Peephole for latency-hiding and hazard-avoidance bet… (PR #191319)
Zeng Wu via llvm-commits
llvm-commits at lists.llvm.org
Fri Apr 17 18:03:56 PDT 2026
https://github.com/zwu-2025 updated https://github.com/llvm/llvm-project/pull/191319
>From 7fd61ff4349809df6db7acd37054ec4a80d41389 Mon Sep 17 00:00:00 2001
From: zwu-2025 <Zeng.Wu2 at amd.com>
Date: Fri, 3 Apr 2026 06:53:01 +0000
Subject: [PATCH 1/5] [AMDGPU] Post-RA Peephole for latency-hiding and
hazard-avoidance between buffer_load and mfma
pattern 1: s_mov_b32 m0 --> s_nop 0 --> buffer_load --> mfma
pattern 2: s_mov_b32 m0 --> buffer_load --> mfma
Result: swap buffer_load and mfma, the remove s_nop 0
---
llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp | 103 ++++++++++++++++++
...gcn-pre-emit-peephole-buffer-load-mfma.mir | 66 +++++++++++
2 files changed, 169 insertions(+)
create mode 100644 llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir
diff --git a/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp b/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
index 411c14fef63b6..25250b932604e 100644
--- a/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
@@ -21,10 +21,15 @@
#include "AMDGPU.h"
#include "GCNSubtarget.h"
#include "MCTargetDesc/AMDGPUMCTargetDesc.h"
+#include "SIRegisterInfo.h"
+#include "Utils/AMDGPUBaseInfo.h"
#include "llvm/ADT/SetVector.h"
+#include "llvm/CodeGen/MachineBasicBlock.h"
#include "llvm/CodeGen/MachineDominators.h"
#include "llvm/CodeGen/MachineFunctionPass.h"
+#include "llvm/CodeGen/MachineInstr.h"
#include "llvm/CodeGen/MachineLoopInfo.h"
+#include "llvm/CodeGen/MachineOperand.h"
#include "llvm/CodeGen/MachinePostDominators.h"
#include "llvm/CodeGen/TargetSchedule.h"
#include "llvm/Support/BranchProbability.h"
@@ -80,6 +85,7 @@ class SIPreEmitPeephole {
// appropriate source modifers and operands into the unpacked instructions.
void addOperandAndMods(MachineInstrBuilder &NewMI, unsigned SrcMods,
bool IsHiBits, const MachineOperand &SrcMO);
+ bool optimizeBufferLoadM0(MachineBasicBlock &MBB);
public:
bool run(MachineFunction &MF, MachineLoopInfo *MLI);
@@ -726,6 +732,97 @@ void SIPreEmitPeephole::performF32Unpacking(MachineInstr &I) {
I.eraseFromParent();
}
+static bool IsVecBufferLoad(const SIInstrInfo *TII, const MachineInstr &MI) {
+ return MI.mayLoad() && (TII->isMTBUF(MI) || TII->isMUBUF(MI));
+};
+
+static MachineInstr *getPrevNonDebugInst(MachineInstr *MI) {
+ for (MachineInstr *I = MI->getPrevNode(); I; I = I->getPrevNode())
+ if (!I->isDebugInstr())
+ return I;
+ return nullptr;
+}
+
+// return true if all of three are true:
+// no flow dependency from BufferLoad to MFMA
+// no anti-flow dependency from BufferLoad to MFMA
+// no output dependency between BufferLoad and MFMA
+static bool canbeReordered(MachineInstr *BufferLoad, MachineInstr *MFMA,
+ const SIRegisterInfo *TRI) {
+ bool HasFlowDep =
+ llvm::any_of(BufferLoad->defs(), [MFMA, TRI](const auto &def) {
+ return def.isReg() && MFMA->readsRegister(def.getReg(), TRI);
+ });
+
+ bool HasAntiDep =
+ llvm::any_of(MFMA->defs(), [BufferLoad, TRI](const auto &def) {
+ return def.isReg() && BufferLoad->readsRegister(def.getReg(), TRI);
+ });
+
+ bool HasOutputDep =
+ llvm::any_of(MFMA->defs(), [BufferLoad, TRI](const auto &def) {
+ return def.isReg() && BufferLoad->modifiesRegister(def.getReg(), TRI);
+ });
+
+ return !HasFlowDep && !HasAntiDep && !HasOutputDep;
+}
+
+// pattern 1: s_mov_b32 m0 --> s_nop 0 --> buffer_load --> mfma
+// pattern 2: s_mov_b32 m0 --> buffer_load --> mfma
+// swap buffer_load and mfma, the remove s_nop 0
+bool SIPreEmitPeephole::optimizeBufferLoadM0(MachineBasicBlock &MBB) {
+ if (MBB.empty()) {
+ return false;
+ }
+
+ bool Changed = false;
+ using InstrIt = MachineBasicBlock::iterator;
+ for (InstrIt I = MBB.begin(), E = MBB.end(); I != E; ++I) {
+ if (!TII->isMFMA(I->getOpcode()))
+ continue;
+ MachineInstr *MFMA = &*I;
+
+ if (I == MBB.begin())
+ continue;
+ MachineInstr *Prev = getPrevNonDebugInst(&*I);
+ if (Prev == nullptr || !IsVecBufferLoad(TII, *Prev) ||
+ !Prev->readsRegister(AMDGPU::M0, TRI))
+ continue;
+
+ MachineInstr *BufferLoad = Prev;
+ Prev = getPrevNonDebugInst(Prev);
+ if (Prev == nullptr)
+ continue;
+
+ MachineInstr *MaybeNop = Prev;
+ MachineInstr *MaybeSMov = nullptr;
+ // In the current lowering pipeline, the post-RA-hazard-rec is just after
+ // this pass, so there is no s_nop. But we still do the check here in case
+ // the order of passes is changed.
+ if (MaybeNop->getOpcode() == AMDGPU::S_NOP &&
+ MaybeNop->getOperand(0).getImm() == 0) {
+ MaybeSMov = getPrevNonDebugInst(MaybeNop);
+ } else {
+ MaybeSMov = MaybeNop;
+ MaybeNop = nullptr;
+ }
+
+ if (MaybeSMov == nullptr || MaybeSMov->getOpcode() != AMDGPU::S_MOV_B32 ||
+ !MaybeSMov->modifiesRegister(AMDGPU::M0, TRI))
+ continue;
+
+ if (!canbeReordered(BufferLoad, MFMA, TRI))
+ continue;
+
+ Changed = true;
+ MBB.splice(std::next(I), &MBB, BufferLoad);
+ if (MaybeNop != nullptr)
+ MaybeNop->eraseFromParent();
+ }
+
+ return Changed;
+}
+
MachineInstrBuilder SIPreEmitPeephole::createUnpackedMI(MachineInstr &I,
uint32_t UnpackedOpcode,
bool IsHiBits) {
@@ -867,5 +964,11 @@ bool SIPreEmitPeephole::run(MachineFunction &MF, MachineLoopInfo *LoopInfo) {
}
}
+ if (ST.hasGFX950Insts()) {
+ for (MachineBasicBlock &MBB : MF) {
+ Changed |= optimizeBufferLoadM0(MBB);
+ }
+ }
+
return Changed;
}
diff --git a/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir b/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir
new file mode 100644
index 0000000000000..005418ddac738
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir
@@ -0,0 +1,66 @@
+# RUN: llc -march=amdgcn -mcpu=gfx950 -run-pass si-pre-emit-peephole -verify-machineinstrs %s -o - | FileCheck %s
+
+---
+name: test_m0_sched_basic
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $sgpr0_sgpr1_sgpr2_sgpr3, $sgpr4, $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5, $vgpr6, $vgpr7, $vgpr8, $vgpr9, $vgpr10, $vgpr11, $vgpr12, $vgpr13, $vgpr14, $vgpr15, $vgpr16, $vgpr17, $vgpr18, $vgpr19, $vgpr20, $vgpr21
+
+ ; CHECK-LABEL: bb.0:
+ ; CHECK: S_MOV_B32
+ ; CHECK: V_MFMA_F32
+ ; CHECK: BUFFER_LOAD_DWORD_IDXEN
+
+ $m0 = S_MOV_B32 $sgpr4
+ S_NOP 0
+ $vgpr32 = BUFFER_LOAD_DWORD_IDXEN $vgpr64, $sgpr0_sgpr1_sgpr2_sgpr3, $sgpr4, 0, 0, 0, implicit $m0, implicit $exec
+ renamable $vgpr0_vgpr1_vgpr2_vgpr3 = nofpexcept V_MFMA_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64 $vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11, $vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19, killed $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, implicit $mode, implicit $exec
+
+
+---
+name: test_m0_sched_basic_no_nop
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $sgpr0_sgpr1_sgpr2_sgpr3, $sgpr4, $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5, $vgpr6, $vgpr7, $vgpr8, $vgpr9, $vgpr10, $vgpr11, $vgpr12, $vgpr13, $vgpr14, $vgpr15, $vgpr16, $vgpr17, $vgpr18, $vgpr19, $vgpr20, $vgpr21
+
+ ; CHECK-LABEL: bb.0:
+ ; CHECK: S_MOV_B32
+ ; CHECK: V_MFMA_F32
+ ; CHECK: BUFFER_LOAD_DWORD_IDXEN
+
+ $m0 = S_MOV_B32 $sgpr4
+ $vgpr32 = BUFFER_LOAD_DWORD_IDXEN $vgpr64, $sgpr0_sgpr1_sgpr2_sgpr3, $sgpr4, 0, 0, 0, implicit $m0, implicit $exec
+ renamable $vgpr0_vgpr1_vgpr2_vgpr3 = nofpexcept V_MFMA_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64 $vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11, $vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19, killed $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, implicit $mode, implicit $exec
+
+---
+name: test_output_dependency
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $sgpr0_sgpr1_sgpr2_sgpr3, $sgpr4, $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5, $vgpr6, $vgpr7, $vgpr8, $vgpr9, $vgpr10, $vgpr11, $vgpr12, $vgpr13, $vgpr14, $vgpr15, $vgpr16, $vgpr17, $vgpr18, $vgpr19, $vgpr20, $vgpr21
+
+ ; CHECK-LABEL: bb.0:
+ ; CHECK: BUFFER_LOAD_DWORD_IDXEN
+ ; CHECK: V_MFMA_F32
+
+ $m0 = S_MOV_B32 $sgpr4
+ $vgpr0 = BUFFER_LOAD_DWORD_IDXEN $vgpr1, $sgpr0_sgpr1_sgpr2_sgpr3, $sgpr4, 0, 0, 0, implicit $m0, implicit $exec
+ renamable $vgpr0_vgpr1_vgpr2_vgpr3 = nofpexcept V_MFMA_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64 $vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11, $vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19, killed $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, implicit $mode, implicit $exec
+
+
+---
+name: test_no_m0
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $sgpr0_sgpr1_sgpr2_sgpr3, $sgpr4, $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5, $vgpr6, $vgpr7, $vgpr8, $vgpr9, $vgpr10, $vgpr11, $vgpr12, $vgpr13, $vgpr14, $vgpr15, $vgpr16, $vgpr17, $vgpr18, $vgpr19, $vgpr20, $vgpr21
+
+ ; CHECK-LABEL: bb.0:
+ ; CHECK: BUFFER_LOAD_DWORD_OFFEN
+ ; CHECK: V_MFMA_F32
+
+ $m0 = S_MOV_B32 $sgpr4
+ $vgpr32 = BUFFER_LOAD_DWORD_OFFEN $vgpr1, $sgpr0_sgpr1_sgpr2_sgpr3, $sgpr4, 0, 0, 0, implicit $exec
+ renamable $vgpr0_vgpr1_vgpr2_vgpr3 = nofpexcept V_MFMA_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64 $vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11, $vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19, killed $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, implicit $mode, implicit $exec
>From b554575d04c9418edfb3130bfc9f8f26bb8f32ff Mon Sep 17 00:00:00 2001
From: Zeng Wu <zengwu13 at amd.com>
Date: Fri, 10 Apr 2026 10:19:38 -0700
Subject: [PATCH 2/5] Apply suggestion from @shiltian
Co-authored-by: Shilei Tian <i at tianshilei.me>
---
llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp b/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
index 25250b932604e..6482c6295e06f 100644
--- a/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
@@ -732,7 +732,7 @@ void SIPreEmitPeephole::performF32Unpacking(MachineInstr &I) {
I.eraseFromParent();
}
-static bool IsVecBufferLoad(const SIInstrInfo *TII, const MachineInstr &MI) {
+static bool isVecBufferLoad(const SIInstrInfo *TII, const MachineInstr &MI) {
return MI.mayLoad() && (TII->isMTBUF(MI) || TII->isMUBUF(MI));
};
>From 4d61a364d5bc8bc0df15fb706887eb3f0d7b2401 Mon Sep 17 00:00:00 2001
From: root <Zeng.Wu2 at amd.com>
Date: Fri, 10 Apr 2026 17:46:07 +0000
Subject: [PATCH 3/5] comments
---
llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp | 28 ++++++++-----------
...gcn-pre-emit-peephole-buffer-load-mfma.mir | 2 +-
2 files changed, 13 insertions(+), 17 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp b/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
index 25250b932604e..9ae4092019bda 100644
--- a/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
@@ -747,21 +747,21 @@ static MachineInstr *getPrevNonDebugInst(MachineInstr *MI) {
// no flow dependency from BufferLoad to MFMA
// no anti-flow dependency from BufferLoad to MFMA
// no output dependency between BufferLoad and MFMA
-static bool canbeReordered(MachineInstr *BufferLoad, MachineInstr *MFMA,
+static bool canBeReordered(MachineInstr *BufferLoad, MachineInstr *MFMA,
const SIRegisterInfo *TRI) {
bool HasFlowDep =
- llvm::any_of(BufferLoad->defs(), [MFMA, TRI](const auto &def) {
+ llvm::any_of(BufferLoad->defs(), [MFMA, TRI](const MachineOperand &def) {
return def.isReg() && MFMA->readsRegister(def.getReg(), TRI);
});
bool HasAntiDep =
- llvm::any_of(MFMA->defs(), [BufferLoad, TRI](const auto &def) {
- return def.isReg() && BufferLoad->readsRegister(def.getReg(), TRI);
+ llvm::any_of(MFMA->defs(), [BufferLoad, TRI](const MachineOperand &def) {
+ return BufferLoad->readsRegister(def.getReg(), TRI);
});
bool HasOutputDep =
- llvm::any_of(MFMA->defs(), [BufferLoad, TRI](const auto &def) {
- return def.isReg() && BufferLoad->modifiesRegister(def.getReg(), TRI);
+ llvm::any_of(MFMA->defs(), [BufferLoad, TRI](const MachineOperand &def) {
+ return BufferLoad->modifiesRegister(def.getReg(), TRI);
});
return !HasFlowDep && !HasAntiDep && !HasOutputDep;
@@ -771,21 +771,18 @@ static bool canbeReordered(MachineInstr *BufferLoad, MachineInstr *MFMA,
// pattern 2: s_mov_b32 m0 --> buffer_load --> mfma
// swap buffer_load and mfma, the remove s_nop 0
bool SIPreEmitPeephole::optimizeBufferLoadM0(MachineBasicBlock &MBB) {
- if (MBB.empty()) {
+ if (MBB.empty())
return false;
- }
bool Changed = false;
using InstrIt = MachineBasicBlock::iterator;
for (InstrIt I = MBB.begin(), E = MBB.end(); I != E; ++I) {
- if (!TII->isMFMA(I->getOpcode()))
+ if (!TII->isMFMA(I->getOpcode()) || I == MBB.begin())
continue;
MachineInstr *MFMA = &*I;
- if (I == MBB.begin())
- continue;
MachineInstr *Prev = getPrevNonDebugInst(&*I);
- if (Prev == nullptr || !IsVecBufferLoad(TII, *Prev) ||
+ if (!Prev || !IsVecBufferLoad(TII, *Prev) ||
!Prev->readsRegister(AMDGPU::M0, TRI))
continue;
@@ -807,11 +804,11 @@ bool SIPreEmitPeephole::optimizeBufferLoadM0(MachineBasicBlock &MBB) {
MaybeNop = nullptr;
}
- if (MaybeSMov == nullptr || MaybeSMov->getOpcode() != AMDGPU::S_MOV_B32 ||
+ if (!MaybeSMov || MaybeSMov->getOpcode() != AMDGPU::S_MOV_B32 ||
!MaybeSMov->modifiesRegister(AMDGPU::M0, TRI))
continue;
- if (!canbeReordered(BufferLoad, MFMA, TRI))
+ if (!canBeReordered(BufferLoad, MFMA, TRI))
continue;
Changed = true;
@@ -965,9 +962,8 @@ bool SIPreEmitPeephole::run(MachineFunction &MF, MachineLoopInfo *LoopInfo) {
}
if (ST.hasGFX950Insts()) {
- for (MachineBasicBlock &MBB : MF) {
+ for (MachineBasicBlock &MBB : MF)
Changed |= optimizeBufferLoadM0(MBB);
- }
}
return Changed;
diff --git a/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir b/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir
index 005418ddac738..47994a77031b7 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir
+++ b/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir
@@ -1,4 +1,4 @@
-# RUN: llc -march=amdgcn -mcpu=gfx950 -run-pass si-pre-emit-peephole -verify-machineinstrs %s -o - | FileCheck %s
+# RUN: llc -march=amdgcn -mcpu=gfx950 -run-pass=si-pre-emit-peephole %s -o - | FileCheck %s
---
name: test_m0_sched_basic
>From 607558f53a0f2e6c628b568be9cd5b5074d9259b Mon Sep 17 00:00:00 2001
From: root <Zeng.Wu2 at amd.com>
Date: Sat, 18 Apr 2026 00:51:37 +0000
Subject: [PATCH 4/5] comments
---
llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp | 43 ++++++++++++++++---
...gcn-pre-emit-peephole-buffer-load-mfma.mir | 7 ++-
2 files changed, 44 insertions(+), 6 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp b/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
index ed7a915cfa51b..d7e8080c7e8f7 100644
--- a/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
@@ -37,6 +37,16 @@ using namespace llvm;
#define DEBUG_TYPE "si-pre-emit-peephole"
+static cl::opt<bool> EnableReorderBufferLoadAndMFMA(
+ "amdgpu-enable-bufferload-and-mfma", cl::Hidden,
+ cl::desc("Enable BufferLoad and MFMA in post-RA peephole."),
+ cl::init(false));
+
+static cl::opt<bool> MFMABufferLoadRatio(
+ "amdgpu-bufferload-mfma-ratio", cl::Hidden,
+ cl::desc("Ratio of MFMA and BufferLoad Ratio to trigger reorder"),
+ cl::init(4));
+
namespace {
class SIPreEmitPeephole {
@@ -767,22 +777,45 @@ static bool canBeReordered(MachineInstr *BufferLoad, MachineInstr *MFMA,
return !HasFlowDep && !HasAntiDep && !HasOutputDep;
}
-// pattern 1: s_mov_b32 m0 --> s_nop 0 --> buffer_load --> mfma
-// pattern 2: s_mov_b32 m0 --> buffer_load --> mfma
-// swap buffer_load and mfma, the remove s_nop 0
+static bool isContinuousMFMA(const SIInstrInfo *TII, MachineInstr *MFMA,
+ unsigned NumMFMA) {
+ unsigned Num = 0;
+ for (MachineInstr *I = MFMA; I; I = I->getNextNode()) {
+ if (TII->isMFMA(I->getOpcode())) {
+ if (++Num == NumMFMA)
+ return true;
+ } else {
+ // There is no continuous MFMA, so false is returned.
+ return false;
+ }
+ }
+ return false;
+}
+
+// pattern 2: s_mov_b32 m0 --> buffer_load --> mfma --> ... -->mfma
+// swap buffer_load and the 1st mfma
bool SIPreEmitPeephole::optimizeBufferLoadM0(MachineBasicBlock &MBB) {
if (MBB.empty())
return false;
bool Changed = false;
using InstrIt = MachineBasicBlock::iterator;
+ unsigned NumberMFMA = MFMABufferLoadRatio;
for (InstrIt I = MBB.begin(), E = MBB.end(); I != E; ++I) {
if (!TII->isMFMA(I->getOpcode()) || I == MBB.begin())
continue;
MachineInstr *MFMA = &*I;
+ // The cycles of buffer_load varies wildly and cycles of MFMA varies
+ // depending on the shape. Option amdgpu-bufferload-mfma-ratio is exposed to
+ // user to set it in perf-tuning if amdgpu-enable-bufferload-and-mfma is
+ // enabled. For example, on GFX950, if mfma is v_mfma_f32_16x16x32_f16, then
+ // ratio with 4 is best to hide latency.
+ if (!isContinuousMFMA(TII, MFMA, NumberMFMA))
+ continue;
+
MachineInstr *Prev = getPrevNonDebugInst(&*I);
- if (!Prev || !IsVecBufferLoad(TII, *Prev) ||
+ if (!Prev || !isVecBufferLoad(TII, *Prev) ||
!Prev->readsRegister(AMDGPU::M0, TRI))
continue;
@@ -961,7 +994,7 @@ bool SIPreEmitPeephole::run(MachineFunction &MF, MachineLoopInfo *LoopInfo) {
}
}
- if (ST.hasGFX950Insts()) {
+ if (ST.hasGFX950Insts() && EnableReorderBufferLoadAndMFMA) {
for (MachineBasicBlock &MBB : MF)
Changed |= optimizeBufferLoadM0(MBB);
}
diff --git a/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir b/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir
index 47994a77031b7..4a8b5b6d04b98 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir
+++ b/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir
@@ -1,4 +1,5 @@
-# RUN: llc -march=amdgcn -mcpu=gfx950 -run-pass=si-pre-emit-peephole %s -o - | FileCheck %s
+# RUN: llc -march=amdgcn -mcpu=gfx950 -run-pass=si-pre-emit-peephole --amdgpu-enable-bufferload-and-mfma=true --amdgpu-bufferload-mfma-ratio=1 %s -o - | FileCheck %s
+# RUN: llc -march=amdgcn -mcpu=gfx950 -run-pass=si-pre-emit-peephole %s -o - | FileCheck %s --check-prefix=DISABLE-REORDER
---
name: test_m0_sched_basic
@@ -12,6 +13,10 @@ body: |
; CHECK: V_MFMA_F32
; CHECK: BUFFER_LOAD_DWORD_IDXEN
+ ; DISABLE-REORDER-LABEL: bb.0:
+ ; DISABLE-REORDER: S_MOV_B32
+ ; DISABLE-REORDER: BUFFER_LOAD_DWORD_IDXEN
+ ; DISABLE-REORDER: V_MFMA_F32
$m0 = S_MOV_B32 $sgpr4
S_NOP 0
$vgpr32 = BUFFER_LOAD_DWORD_IDXEN $vgpr64, $sgpr0_sgpr1_sgpr2_sgpr3, $sgpr4, 0, 0, 0, implicit $m0, implicit $exec
>From 5a84a4f2a090d564662a63e3621b9f99e19ce330 Mon Sep 17 00:00:00 2001
From: root <Zeng.Wu2 at amd.com>
Date: Sat, 18 Apr 2026 01:03:39 +0000
Subject: [PATCH 5/5] rename
---
llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp | 11 ++++++-----
.../amdgcn-pre-emit-peephole-buffer-load-mfma.mir | 2 +-
2 files changed, 7 insertions(+), 6 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp b/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
index d7e8080c7e8f7..bb0738d229922 100644
--- a/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
@@ -38,8 +38,8 @@ using namespace llvm;
#define DEBUG_TYPE "si-pre-emit-peephole"
static cl::opt<bool> EnableReorderBufferLoadAndMFMA(
- "amdgpu-enable-bufferload-and-mfma", cl::Hidden,
- cl::desc("Enable BufferLoad and MFMA in post-RA peephole."),
+ "amdgpu-enable-reorder-bufferload-and-mfma", cl::Hidden,
+ cl::desc("Enable Reorder BufferLoad and MFMA in post-RA peephole."),
cl::init(false));
static cl::opt<bool> MFMABufferLoadRatio(
@@ -808,9 +808,10 @@ bool SIPreEmitPeephole::optimizeBufferLoadM0(MachineBasicBlock &MBB) {
// The cycles of buffer_load varies wildly and cycles of MFMA varies
// depending on the shape. Option amdgpu-bufferload-mfma-ratio is exposed to
- // user to set it in perf-tuning if amdgpu-enable-bufferload-and-mfma is
- // enabled. For example, on GFX950, if mfma is v_mfma_f32_16x16x32_f16, then
- // ratio with 4 is best to hide latency.
+ // user to set it in perf-tuning if
+ // amdgpu-enable-reorder-bufferload-and-mfma is enabled. For example, on
+ // GFX950, if mfma is v_mfma_f32_16x16x32_f16, then ratio with 4 is best to
+ // hide latency.
if (!isContinuousMFMA(TII, MFMA, NumberMFMA))
continue;
diff --git a/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir b/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir
index 4a8b5b6d04b98..c5aaf823d2da3 100644
--- a/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir
+++ b/llvm/test/CodeGen/AMDGPU/amdgcn-pre-emit-peephole-buffer-load-mfma.mir
@@ -1,4 +1,4 @@
-# RUN: llc -march=amdgcn -mcpu=gfx950 -run-pass=si-pre-emit-peephole --amdgpu-enable-bufferload-and-mfma=true --amdgpu-bufferload-mfma-ratio=1 %s -o - | FileCheck %s
+# RUN: llc -march=amdgcn -mcpu=gfx950 -run-pass=si-pre-emit-peephole --amdgpu-enable-reorder-bufferload-and-mfma=true --amdgpu-bufferload-mfma-ratio=1 %s -o - | FileCheck %s
# RUN: llc -march=amdgcn -mcpu=gfx950 -run-pass=si-pre-emit-peephole %s -o - | FileCheck %s --check-prefix=DISABLE-REORDER
---
More information about the llvm-commits
mailing list