[llvm] [AMDGPU] Introduce GCNPostGenericScheduler to model MFMA-VALU coexecution (PR #184084)
Anshil Gandhi via llvm-commits
llvm-commits at lists.llvm.org
Thu Apr 16 11:49:11 PDT 2026
https://github.com/gandhi56 updated https://github.com/llvm/llvm-project/pull/184084
>From 2ff8a55439a143fb05e9140115892465ef109a96 Mon Sep 17 00:00:00 2001
From: Anshil Gandhi <Anshil.Gandhi at amd.com>
Date: Thu, 12 Mar 2026 10:45:21 -0500
Subject: [PATCH 1/2] [AMDGPU] Introduce GCNPostGenericScheduler to model
MFMA-VALU nop hazard in post-RA scheduling
The AMDGPU hazard recognizer inserts a nop between
an MFMA and a subsequent non-MFMA VALU instruction.
The generic PostGenericScheduler is unaware of this
penalty, so it may place a VALU directly after an
MFMA, resulting in avoidable nop stalls.
This patch introduces GCNPostGenericScheduler, a
target-specific subclass of PostGenericScheduler,
and wires it into the post-RA machine scheduler.
Its schedNode override tracks whether the last
top-scheduled instruction was an MFMA; when a
non-MFMA VALU follows, it bumps the top cycle
boundary by one to reflect the nop cost. This
gives the scheduler the information it needs to
fill the hazard slot with an independent
instruction rather than leaving a dead cycle.
The test diffs show MFMA instructions being pulled
earlier relative to non-MFMA VALUs across the iglp
and rewrite-vgpr-mfma-to-agpr tests, with several
s_nop counts reduced (e.g. 6 -> 3, 3 -> 2).
---
.../lib/Target/AMDGPU/AMDGPUTargetMachine.cpp | 6 +-
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 36 ++++++++
llvm/lib/Target/AMDGPU/GCNSchedStrategy.h | 20 +++++
.../AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir | 86 +++++++++----------
.../AMDGPU/llvm.amdgcn.iglp.opt.exp.small.mir | 6 +-
.../CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.ll | 2 +-
.../AMDGPU/rewrite-vgpr-mfma-to-agpr.ll | 16 ++--
7 files changed, 114 insertions(+), 58 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
index d7e2264e51fd6..29f6624e89b0f 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -1319,9 +1319,9 @@ GCNTargetMachine::createPostMachineScheduler(MachineSchedContext *C) const {
if (useNoopPostScheduler(C->MF->getFunction()))
return createGCNNoopPostMachineScheduler(C);
- ScheduleDAGMI *DAG =
- new GCNPostScheduleDAGMILive(C, std::make_unique<PostGenericScheduler>(C),
- /*RemoveKillFlags=*/true);
+ ScheduleDAGMI *DAG = new GCNPostScheduleDAGMILive(
+ C, std::make_unique<GCNPostGenericScheduler>(C),
+ /*RemoveKillFlags=*/true);
const GCNSubtarget &ST = C->MF->getSubtarget<GCNSubtarget>();
DAG->addMutation(createLoadClusterDAGMutation(DAG->TII, DAG->TRI));
if (ST.shouldClusterStores())
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index af2b2188c3081..54d0dedd8e78d 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -3274,6 +3274,42 @@ static bool hasIGLPInstrs(ScheduleDAGInstrs *DAG) {
});
}
+void GCNPostGenericScheduler::initPolicy(MachineBasicBlock::iterator Begin,
+ MachineBasicBlock::iterator End,
+ unsigned NumRegionInstrs) {
+ LastTopScheduledIsMFMA = false;
+ LastBottomScheduledIsMFMA = false;
+ PostGenericScheduler::initPolicy(Begin, End, NumRegionInstrs);
+}
+
+void GCNPostGenericScheduler::schedNode(SUnit *SU, bool IsTopNode) {
+ // Track MFMA vs other VALU along each scheduling direction. Issuing a
+ // non-MFMA VALU immediately after an MFMA is modeled as costing an extra
+ // cycle (bumpCycle) so the post-scheduler's pressure/latency accounting
+ // matches hardware behavior for that transition.
+ if (SU->isInstr()) {
+ MachineInstr *MI = SU->getInstr();
+ bool IsNonMFMAVALU = SIInstrInfo::isVALU(*MI) && !SIInstrInfo::isMFMA(*MI);
+ if (IsNonMFMAVALU) {
+ if (IsTopNode && LastTopScheduledIsMFMA) {
+ Top.bumpCycle(Top.getCurrCycle() + 1);
+ LastTopScheduledIsMFMA = false;
+ }
+ if (!IsTopNode && LastBottomScheduledIsMFMA) {
+ Bot.bumpCycle(Bot.getCurrCycle() + 1);
+ LastBottomScheduledIsMFMA = false;
+ }
+ } else if (SIInstrInfo::isMFMA(*MI)) {
+ if (IsTopNode)
+ LastTopScheduledIsMFMA = true;
+ if (!IsTopNode)
+ LastBottomScheduledIsMFMA = true;
+ }
+ }
+
+ PostGenericScheduler::schedNode(SU, IsTopNode);
+}
+
GCNPostScheduleDAGMILive::GCNPostScheduleDAGMILive(
MachineSchedContext *C, std::unique_ptr<MachineSchedStrategy> S,
bool RemoveKillFlags)
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
index 2cc9e81a65191..5b0ab89717401 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
@@ -808,6 +808,26 @@ class MemoryClauseInitialScheduleStage : public GCNSchedStage {
: GCNSchedStage(StageID, DAG) {}
};
+/// AMDGPU post-RA scheduler built on \c PostGenericScheduler. Extends the
+/// generic policy with VALU/MFMA awareness: when a non-MFMA VALU is placed
+/// immediately after an MFMA on the same scheduling boundary, the boundary
+/// cycle is advanced to reflect the hardware cost of that issue pattern.
+class GCNPostGenericScheduler : public PostGenericScheduler {
+private:
+ bool LastTopScheduledIsMFMA = false;
+ bool LastBottomScheduledIsMFMA = false;
+
+public:
+ GCNPostGenericScheduler(const MachineSchedContext *C)
+ : PostGenericScheduler(C) {}
+
+ void initPolicy(MachineBasicBlock::iterator Begin,
+ MachineBasicBlock::iterator End,
+ unsigned NumRegionInstrs) override;
+
+ void schedNode(SUnit *SU, bool IsTopNode) override;
+};
+
class GCNPostScheduleDAGMILive final : public ScheduleDAGMI {
private:
std::vector<std::unique_ptr<ScheduleDAGMutation>> SavedMutations;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir
index 94de6dd31cad5..d3e6ae7482e03 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir
@@ -240,8 +240,8 @@
; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: v_perm_b32 v239, v174, v172, s5
; GCN-NEXT: v_perm_b32 v241, v174, v172, s7
- ; GCN-NEXT: v_perm_b32 v243, v175, v173, s5
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[214:215], v[144:145], v[64:79]
+ ; GCN-NEXT: v_perm_b32 v243, v175, v173, s5
; GCN-NEXT: v_perm_b32 v245, v175, v173, s7
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[176:177], v[156:157], v[96:111]
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[220:221], v[142:143], v[112:127]
@@ -362,8 +362,8 @@
; GCN-NEXT: v_perm_b32 v192, v195, v193, s5
; GCN-NEXT: v_perm_b32 v194, v195, v193, s7
; GCN-NEXT: v_perm_b32 v193, v221, v219, s5
- ; GCN-NEXT: v_perm_b32 v195, v221, v219, s7
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[166:167], v[146:147], v[80:95]
+ ; GCN-NEXT: v_perm_b32 v195, v221, v219, s7
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[168:169], v[140:141], v[80:95]
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[170:171], v[142:143], v[80:95]
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[172:173], v[148:149], v[80:95]
@@ -387,61 +387,61 @@
; GCN-NEXT: v_mul_f32_e32 v218, s4, v114
; GCN-NEXT: v_mul_f32_e32 v219, s4, v115
; GCN-NEXT: v_max3_f32 v213, v213, v218, v219
- ; GCN-NEXT: v_mul_f32_e32 v218, s4, v116
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[182:183], v[134:135], v[80:95]
+ ; GCN-NEXT: v_mul_f32_e32 v218, s4, v116
; GCN-NEXT: v_mul_f32_e32 v219, s4, v117
; GCN-NEXT: v_max3_f32 v213, v213, v218, v219
; GCN-NEXT: v_mul_f32_e32 v218, s4, v118
; GCN-NEXT: v_mul_f32_e32 v219, s4, v119
; GCN-NEXT: v_max3_f32 v213, v213, v218, v219
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[152:153], v[132:133], v[64:79]
; GCN-NEXT: v_mul_f32_e32 v218, s4, v120
; GCN-NEXT: v_mul_f32_e32 v219, s4, v121
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[152:153], v[132:133], v[64:79]
; GCN-NEXT: v_max3_f32 v213, v213, v218, v219
; GCN-NEXT: v_mul_f32_e32 v218, s4, v122
; GCN-NEXT: v_mul_f32_e32 v219, s4, v123
; GCN-NEXT: v_max3_f32 v213, v213, v218, v219
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[236:237], v[130:131], v[96:111]
; GCN-NEXT: v_mul_f32_e32 v218, s4, v124
; GCN-NEXT: v_mul_f32_e32 v219, s4, v125
; GCN-NEXT: v_max3_f32 v213, v213, v218, v219
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[236:237], v[130:131], v[96:111]
; GCN-NEXT: v_mul_f32_e32 v218, s4, v126
; GCN-NEXT: v_mul_f32_e32 v219, s4, v127
; GCN-NEXT: v_max3_f32 v213, v213, v218, v219
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[184:185], v[128:129], v[80:95]
- ; GCN-NEXT: s_nop 6
+ ; GCN-NEXT: s_nop 3
; GCN-NEXT: v_mul_f32_e32 v214, s4, v96
; GCN-NEXT: v_mul_f32_e32 v215, s4, v97
; GCN-NEXT: v_max3_f32 v213, v213, v214, v215
; GCN-NEXT: v_mul_f32_e32 v214, s4, v98
; GCN-NEXT: v_mul_f32_e32 v215, s4, v99
; GCN-NEXT: v_max3_f32 v213, v213, v214, v215
- ; GCN-NEXT: v_mul_f32_e32 v214, s4, v100
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[154:155], v[134:135], v[64:79]
+ ; GCN-NEXT: v_mul_f32_e32 v214, s4, v100
; GCN-NEXT: v_mul_f32_e32 v215, s4, v101
; GCN-NEXT: v_max3_f32 v213, v213, v214, v215
; GCN-NEXT: v_mul_f32_e32 v214, s4, v102
; GCN-NEXT: v_mul_f32_e32 v215, s4, v103
; GCN-NEXT: v_max3_f32 v213, v213, v214, v215
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[186:187], v[130:131], v[80:95]
; GCN-NEXT: v_mul_f32_e32 v214, s4, v104
; GCN-NEXT: v_mul_f32_e32 v215, s4, v105
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[186:187], v[130:131], v[80:95]
; GCN-NEXT: v_max3_f32 v213, v213, v214, v215
; GCN-NEXT: v_mul_f32_e32 v214, s4, v106
; GCN-NEXT: v_mul_f32_e32 v215, s4, v107
; GCN-NEXT: v_max3_f32 v213, v213, v214, v215
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[156:157], v[128:129], v[64:79]
; GCN-NEXT: v_mul_f32_e32 v214, s4, v108
; GCN-NEXT: v_mul_f32_e32 v215, s4, v109
; GCN-NEXT: v_max3_f32 v213, v213, v214, v215
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[156:157], v[128:129], v[64:79]
; GCN-NEXT: v_mul_f32_e32 v214, s4, v110
; GCN-NEXT: v_mul_f32_e32 v215, s4, v111
; GCN-NEXT: v_max3_f32 v213, v213, v214, v215
; GCN-NEXT: v_mul_f32_e32 v140, s4, v80
; GCN-NEXT: v_mul_f32_e32 v141, s4, v81
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[158:159], v[130:131], v[64:79]
; GCN-NEXT: v_max3_f32 v140, v213, v140, v141
; GCN-NEXT: v_mul_f32_e32 v141, s4, v82
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[158:159], v[130:131], v[64:79]
; GCN-NEXT: v_mul_f32_e32 v142, s4, v83
; GCN-NEXT: v_max3_f32 v140, v140, v141, v142
; GCN-NEXT: v_mul_f32_e32 v141, s4, v84
@@ -602,15 +602,13 @@
; GCN-NEXT: v_cvt_f16_f32_e32 v126, v121
; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v125
; GCN-NEXT: v_fma_f32 v139, s4, v96, -v128
- ; GCN-NEXT: v_fma_f32 v127, s4, v127, -v128
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[142:143], v[146:147], v[48:63]
; GCN-NEXT: v_exp_f32_e32 v123, v150
+ ; GCN-NEXT: v_fma_f32 v127, s4, v127, -v128
; GCN-NEXT: v_mul_f32_e32 v127, 0x3fb8aa3b, v127
; GCN-NEXT: v_fma_f32 v143, s4, v101, -v128
; GCN-NEXT: v_fma_f32 v64, s4, v64, -v128
; GCN-NEXT: v_fma_f32 v65, s4, v65, -v128
- ; GCN-NEXT: v_fma_f32 v68, s4, v68, -v128
- ; GCN-NEXT: v_fma_f32 v69, s4, v69, -v128
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[134:135], v[0:15]
; GCN-NEXT: v_exp_f32_e32 v124, v151
; GCN-NEXT: ds_read_b128 v[130:133], v197
@@ -619,6 +617,8 @@
; GCN-NEXT: ds_read_b128 v[146:149], v197 offset:576
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: v_fma_f32 v68, s4, v68, -v128
+ ; GCN-NEXT: v_fma_f32 v69, s4, v69, -v128
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[136:137], v[134:135], v[32:47]
; GCN-NEXT: v_cvt_f16_f32_e32 v136, v122
; GCN-NEXT: v_exp_f32_e32 v96, v129
@@ -687,53 +687,53 @@
; GCN-NEXT: buffer_load_dwordx2 v[144:145], v209, s[0:3], 0 offen sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: ;;#ASMSTART
- ; GCN-NEXT: s_waitcnt vmcnt(8)
- ; GCN-NEXT: ;;#ASMEND
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[138:139], v[126:127], v[48:63]
; GCN-NEXT: v_cvt_f16_f32_e32 v126, v99
; GCN-NEXT: v_fma_f32 v127, s4, v103, -v128
; GCN-NEXT: v_exp_f32_e32 v103, v150
- ; GCN-NEXT: v_fma_f32 v139, s4, v105, -v128
+ ; GCN-NEXT: ;;#ASMSTART
+ ; GCN-NEXT: s_waitcnt vmcnt(8)
+ ; GCN-NEXT: ;;#ASMEND
; GCN-NEXT: v_pack_b32_f16 v147, v147, v126
+ ; GCN-NEXT: v_fma_f32 v139, s4, v105, -v128
; GCN-NEXT: v_mul_f32_e32 v138, 0x3fb8aa3b, v127
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[146:147], v[0:15]
; GCN-NEXT: v_perm_b32 v152, v135, v131, s5
; GCN-NEXT: v_perm_b32 v154, v135, v131, s7
; GCN-NEXT: v_fma_f32 v135, s4, v104, -v128
; GCN-NEXT: v_perm_b32 v126, v134, v130, s5
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[146:147], v[0:15]
; GCN-NEXT: v_perm_b32 v150, v134, v130, s7
; GCN-NEXT: v_cvt_f16_f32_e32 v134, v100
; GCN-NEXT: v_exp_f32_e32 v104, v129
; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v135
; GCN-NEXT: v_cvt_f16_f32_e32 v135, v101
- ; GCN-NEXT: ds_read_b128 v[130:133], v198
- ; GCN-NEXT: s_waitcnt lgkmcnt(0)
- ; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_perm_b32 v127, v144, v142, s5
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[148:149], v[146:147], v[32:47]
+ ; GCN-NEXT: v_exp_f32_e32 v105, v125
; GCN-NEXT: v_pack_b32_f16 v148, v134, v135
; GCN-NEXT: v_fma_f32 v135, s4, v106, -v128
- ; GCN-NEXT: v_exp_f32_e32 v105, v125
; GCN-NEXT: v_cvt_f16_f32_e32 v134, v102
- ; GCN-NEXT: v_perm_b32 v151, v144, v142, s7
- ; GCN-NEXT: v_perm_b32 v153, v145, v143, s5
- ; GCN-NEXT: v_perm_b32 v155, v145, v143, s7
+ ; GCN-NEXT: ds_read_b128 v[130:133], v198
+ ; GCN-NEXT: s_waitcnt lgkmcnt(0)
+ ; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: v_perm_b32 v127, v144, v142, s5
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[136:137], v[146:147], v[16:31]
; GCN-NEXT: v_exp_f32_e32 v106, v156
; GCN-NEXT: v_mul_f32_e32 v156, 0x3fb8aa3b, v135
; GCN-NEXT: v_cvt_f16_f32_e32 v135, v103
; GCN-NEXT: v_fma_f32 v136, s4, v107, -v128
+ ; GCN-NEXT: v_perm_b32 v151, v144, v142, s7
+ ; GCN-NEXT: v_perm_b32 v153, v145, v143, s5
+ ; GCN-NEXT: v_perm_b32 v155, v145, v143, s7
; GCN-NEXT: ds_read_b128 v[142:145], v198 offset:576
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v139
- ; GCN-NEXT: v_pack_b32_f16 v149, v134, v135
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[140:141], v[146:147], v[48:63]
; GCN-NEXT: v_mul_f32_e32 v146, 0x3fb8aa3b, v136
+ ; GCN-NEXT: v_pack_b32_f16 v149, v134, v135
; GCN-NEXT: ds_read_b128 v[134:137], v198 offset:1152
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v139
; GCN-NEXT: v_exp_f32_e32 v107, v138
; GCN-NEXT: ds_read_b128 v[138:141], v198 offset:1728
; GCN-NEXT: s_waitcnt lgkmcnt(0)
@@ -784,9 +784,9 @@
; GCN-NEXT: v_cvt_f16_f32_e32 v134, v111
; GCN-NEXT: v_mul_f32_e32 v156, 0x3fb8aa3b, v137
; GCN-NEXT: v_fma_f32 v137, s4, v83, -v128
- ; GCN-NEXT: v_mul_f32_e32 v157, 0x3fb8aa3b, v137
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[140:141], v[142:143], v[48:63]
; GCN-NEXT: v_exp_f32_e32 v83, v135
+ ; GCN-NEXT: v_mul_f32_e32 v157, 0x3fb8aa3b, v137
; GCN-NEXT: v_pack_b32_f16 v145, v136, v134
; GCN-NEXT: ds_read_b128 v[134:137], v197 offset:1152
; GCN-NEXT: s_waitcnt lgkmcnt(0)
@@ -797,13 +797,13 @@
; GCN-NEXT: ;;#ASMSTART
; GCN-NEXT: s_waitcnt vmcnt(8)
; GCN-NEXT: ;;#ASMEND
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[130:131], v[144:145], v[0:15]
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: ds_write_b64 v199, v[126:127]
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GCN-NEXT: ds_write_b64 v200, v[150:151]
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[130:131], v[144:145], v[0:15]
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GCN-NEXT: ds_write_b64 v201, v[152:153]
@@ -812,18 +812,16 @@
; GCN-NEXT: ds_write_b64 v202, v[154:155]
; GCN-NEXT: v_fma_f32 v127, s4, v84, -v128
; GCN-NEXT: v_exp_f32_e32 v84, v129
- ; GCN-NEXT: v_fma_f32 v130, s4, v85, -v128
- ; GCN-NEXT: v_cvt_f16_f32_e32 v126, v80
- ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v127
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[146:147], v[144:145], v[32:47]
+ ; GCN-NEXT: v_fma_f32 v130, s4, v85, -v128
; GCN-NEXT: v_exp_f32_e32 v85, v125
; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v130
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_load_dwordx2 v[130:131], v206, s[0:3], 0 offen sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_cvt_f16_f32_e32 v127, v81
- ; GCN-NEXT: v_pack_b32_f16 v126, v126, v127
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v126, v80
+ ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v127
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[134:135], v[144:145], v[16:31]
; GCN-NEXT: v_fma_f32 v134, s4, v86, -v128
; GCN-NEXT: v_mul_f32_e32 v158, 0x3fb8aa3b, v134
@@ -836,10 +834,12 @@
; GCN-NEXT: buffer_load_dwordx2 v[146:147], v205, s[0:3], 0 offen sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_cvt_f16_f32_e32 v127, v82
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v127, v81
; GCN-NEXT: v_exp_f32_e32 v86, v156
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[138:139], v[144:145], v[48:63]
; GCN-NEXT: v_cvt_f16_f32_e32 v138, v83
+ ; GCN-NEXT: v_pack_b32_f16 v126, v126, v127
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v127, v82
; GCN-NEXT: ;;#ASMSTART
; GCN-NEXT: s_waitcnt vmcnt(8)
; GCN-NEXT: ;;#ASMEND
@@ -898,17 +898,17 @@
; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v130
; GCN-NEXT: v_cvt_f16_f32_e32 v130, v89
; GCN-NEXT: v_fma_f32 v131, s4, v93, -v128
- ; GCN-NEXT: v_pack_b32_f16 v130, v126, v130
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[142:143], v[146:147], v[32:47]
; GCN-NEXT: v_exp_f32_e32 v93, v125
+ ; GCN-NEXT: v_pack_b32_f16 v130, v126, v130
; GCN-NEXT: v_fma_f32 v126, s4, v94, -v128
; GCN-NEXT: v_cvt_f16_f32_e32 v125, v90
; GCN-NEXT: v_mul_f32_e32 v143, 0x3fb8aa3b, v126
; GCN-NEXT: v_cvt_f16_f32_e32 v126, v91
; GCN-NEXT: v_mul_f32_e32 v142, 0x3fb8aa3b, v131
- ; GCN-NEXT: v_fma_f32 v131, s4, v95, -v128
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[134:135], v[146:147], v[16:31]
; GCN-NEXT: v_exp_f32_e32 v94, v148
+ ; GCN-NEXT: v_fma_f32 v131, s4, v95, -v128
; GCN-NEXT: v_cvt_f16_f32_e32 v64, v93
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[138:139], v[146:147], v[48:63]
; GCN-NEXT: v_exp_f32_e32 v95, v127
@@ -951,14 +951,14 @@
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: ds_write_b64 v199, v[150:151]
- ; GCN-NEXT: buffer_wbl2 sc0 sc1
- ; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
- ; GCN-NEXT: ds_write_b64 v200, v[152:153]
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[142:143], v[0:15]
; GCN-NEXT: v_cvt_f16_f32_e32 v132, v125
; GCN-NEXT: v_exp_f32_e32 v130, v158
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
+ ; GCN-NEXT: ds_write_b64 v200, v[152:153]
+ ; GCN-NEXT: buffer_wbl2 sc0 sc1
+ ; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GCN-NEXT: ds_write_b64 v201, v[154:155]
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
@@ -973,14 +973,13 @@
; GCN-NEXT: v_mul_f32_e32 v144, 0x3fb8aa3b, v69
; GCN-NEXT: v_fma_f32 v69, s4, v71, -v128
; GCN-NEXT: v_pack_b32_f16 v140, v132, v68
- ; GCN-NEXT: v_cvt_f16_f32_e32 v68, v129
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[64:65], v[142:143], v[16:31]
; GCN-NEXT: v_cvt_f16_f32_e32 v64, v127
; GCN-NEXT: v_exp_f32_e32 v132, v145
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v68, v129
; GCN-NEXT: v_fma_f32 v65, s4, v70, -v128
; GCN-NEXT: v_mul_f32_e32 v65, 0x3fb8aa3b, v65
; GCN-NEXT: v_fma_f32 v145, s4, v73, -v128
- ; GCN-NEXT: v_mul_f32_e32 v147, 0x3fb8aa3b, v145
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[136:137], v[142:143], v[48:63]
; GCN-NEXT: v_exp_f32_e32 v133, v141
; GCN-NEXT: v_mul_f32_e32 v142, 0x3fb8aa3b, v69
@@ -998,6 +997,7 @@
; GCN-NEXT: ds_read_b128 v[134:137], v198 offset:576
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: v_mul_f32_e32 v147, 0x3fb8aa3b, v145
; GCN-NEXT: v_pack_b32_f16 v64, v64, v143
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[148:149], v[140:141], v[32:47]
; GCN-NEXT: v_exp_f32_e32 v73, v144
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.small.mir b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.small.mir
index 0a8d7acd187fc..77684d1821fe7 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.small.mir
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.small.mir
@@ -148,13 +148,13 @@
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GCN-NEXT: ds_write_b32 v78, v72
- ; GCN-NEXT: v_mul_f32_e32 v74, s4, v20
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[68:69], v[64:65], v[0:15]
; GCN-NEXT: v_mul_f32_e32 v64, s4, v16
; GCN-NEXT: v_mul_f32_e32 v65, s4, v17
; GCN-NEXT: v_mul_f32_e32 v68, s4, v18
; GCN-NEXT: v_mul_f32_e32 v69, s4, v19
; GCN-NEXT: v_max3_f32 v64, v64, s5, v65
+ ; GCN-NEXT: v_mul_f32_e32 v74, s4, v20
; GCN-NEXT: v_mul_f32_e32 v80, s4, v21
; GCN-NEXT: v_max3_f32 v64, v64, v68, v69
; GCN-NEXT: v_mul_f32_e32 v84, s4, v22
@@ -338,8 +338,8 @@
; GCN-NEXT: v_exp_f32_e32 v22, v64
; GCN-NEXT: v_cvt_f16_f32_e32 v21, v88
; GCN-NEXT: v_exp_f32_e32 v64, v65
- ; GCN-NEXT: v_mul_f32_e32 v23, 0x3fb8aa3b, v23
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[24:25], v[18:19], v[32:47]
+ ; GCN-NEXT: v_mul_f32_e32 v23, 0x3fb8aa3b, v23
; GCN-NEXT: v_add_f32_e32 v17, v75, v17
; GCN-NEXT: v_cvt_f16_f32_e32 v18, v30
; GCN-NEXT: v_fma_f32 v24, s4, v3, -v72
@@ -389,8 +389,8 @@
; GCN-NEXT: v_cvt_f16_f32_e32 v29, v65
; GCN-NEXT: v_fma_f32 v10, s4, v10, -v72
; GCN-NEXT: v_exp_f32_e32 v67, v67
- ; GCN-NEXT: v_mul_f32_e32 v6, 0x3fb8aa3b, v6
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[70:71], v[4:5], v[32:47]
+ ; GCN-NEXT: v_mul_f32_e32 v6, 0x3fb8aa3b, v6
; GCN-NEXT: v_add_f32_e32 v17, v83, v17
; GCN-NEXT: v_cvt_f16_f32_e32 v5, v68
; GCN-NEXT: v_exp_f32_e32 v6, v6
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.ll
index a10c99070d8e1..0c96634c99122 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.ll
@@ -86,7 +86,6 @@ define amdgpu_kernel void @test_iglp_opt_mfma_gemm(ptr addrspace(3) noalias %in,
; GCN-NEXT: v_mfma_f32_32x32x1f32 a[64:95], v2, v3, a[64:95]
; GCN-NEXT: ds_write_b128 v0, a[56:59] offset:24672
; GCN-NEXT: ds_write_b128 v0, a[60:63] offset:24688
-; GCN-NEXT: ds_write_b128 v0, a[48:51] offset:24640
; GCN-NEXT: ds_write_b128 v0, a[120:123] offset:8288
; GCN-NEXT: ds_write_b128 v0, a[124:127] offset:8304
; GCN-NEXT: ds_write_b128 v0, a[112:115] offset:8256
@@ -95,6 +94,7 @@ define amdgpu_kernel void @test_iglp_opt_mfma_gemm(ptr addrspace(3) noalias %in,
; GCN-NEXT: ds_write_b128 v0, a[108:111] offset:8240
; GCN-NEXT: ds_write_b128 v0, a[96:99] offset:8192
; GCN-NEXT: ds_write_b128 v0, a[100:103] offset:8208
+; GCN-NEXT: ds_write_b128 v0, a[48:51] offset:24640
; GCN-NEXT: ds_write_b128 v0, a[52:55] offset:24656
; GCN-NEXT: ds_write_b128 v0, a[40:43] offset:24608
; GCN-NEXT: ds_write_b128 v0, a[44:47] offset:24624
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll
index 5aa2cc713a4ec..dddc9182677f8 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll
@@ -392,15 +392,15 @@ define amdgpu_kernel void @illegal_mfma_after_rewrite() #1 {
; CHECK-NEXT: v_mfma_f32_16x16x16_f16 v[4:7], v[26:27], v[26:27], v[0:3]
; CHECK-NEXT: v_mov_b64_e32 v[28:29], s[0:1]
; CHECK-NEXT: s_mov_b32 s0, 0x7e007e00
-; CHECK-NEXT: s_mov_b32 s1, s0
; CHECK-NEXT: v_mfma_f32_16x16x16_f16 v[4:7], v[26:27], v[26:27], v[4:7]
+; CHECK-NEXT: s_mov_b32 s1, s0
; CHECK-NEXT: v_mov_b64_e32 v[30:31], s[0:1]
; CHECK-NEXT: v_mfma_f32_16x16x16_f16 v[6:9], v[26:27], v[28:29], v[0:3]
; CHECK-NEXT: v_mfma_f32_16x16x16_f16 v[6:9], v[26:27], v[26:27], v[6:9]
-; CHECK-NEXT: s_nop 3
+; CHECK-NEXT: s_nop 2
; CHECK-NEXT: v_cvt_f16_f32_e32 v24, v4
; CHECK-NEXT: v_mfma_f32_16x16x16_f16 v[12:15], v[26:27], v[30:31], v[0:3]
-; CHECK-NEXT: s_nop 0
+; CHECK-NEXT: s_nop 1
; CHECK-NEXT: v_mov_b32_e32 v8, 0x7fc00000
; CHECK-NEXT: v_mov_b32_e32 v9, v8
; CHECK-NEXT: v_mov_b32_e32 v10, v8
@@ -425,8 +425,9 @@ define amdgpu_kernel void @illegal_mfma_after_rewrite() #1 {
; CHECK-NEXT: buffer_inv sc0 sc1
; CHECK-NEXT: s_nop 1
; CHECK-NEXT: v_cvt_f16_f32_e32 v6, v6
-; CHECK-NEXT: global_store_short v[0:1], v6, off
; CHECK-NEXT: v_mfma_f32_16x16x16_f16 v[16:19], v[26:27], v[26:27], v[20:23]
+; CHECK-NEXT: global_store_short v[0:1], v6, off
+; CHECK-NEXT: v_mfma_f32_16x16x16_f16 v[2:5], v[28:29], v[26:27], v[2:5]
; CHECK-NEXT: buffer_wbl2 sc0 sc1
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: buffer_inv sc0 sc1
@@ -434,8 +435,6 @@ define amdgpu_kernel void @illegal_mfma_after_rewrite() #1 {
; CHECK-NEXT: buffer_wbl2 sc0 sc1
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: buffer_inv sc0 sc1
-; CHECK-NEXT: v_mfma_f32_16x16x16_f16 v[2:5], v[28:29], v[26:27], v[2:5]
-; CHECK-NEXT: s_nop 6
; CHECK-NEXT: v_cvt_f16_f32_e32 v6, v2
; CHECK-NEXT: v_mfma_f32_16x16x16_f16 v[2:5], v[30:31], v[26:27], v[16:19]
; CHECK-NEXT: global_store_short v[0:1], v6, off
@@ -723,9 +722,10 @@ define void @test_rewrite_mfma_copy_from_agpr_class_f64_4x4x4f64_chain(double %a
; CHECK-NEXT: v_and_b32_e32 v2, 0x3ff, v31
; CHECK-NEXT: v_lshlrev_b32_e32 v2, 3, v2
; CHECK-NEXT: v_mov_b32_e32 v3, 0
-; CHECK-NEXT: v_lshl_add_u64 v[2:3], v[8:9], 0, v[2:3]
+; CHECK-NEXT: s_nop 0
; CHECK-NEXT: v_mfma_f64_4x4x4_4b_f64 a[0:1], v[4:5], v[6:7], a[0:1]
-; CHECK-NEXT: s_nop 8
+; CHECK-NEXT: v_lshl_add_u64 v[2:3], v[8:9], 0, v[2:3]
+; CHECK-NEXT: s_nop 7
; CHECK-NEXT: global_store_dwordx2 v[2:3], a[0:1], off
; CHECK-NEXT: s_waitcnt vmcnt(0)
; CHECK-NEXT: s_setpc_b64 s[30:31]
>From 237a1d9c71c35303a6e8a6419330fe124d00e7bf Mon Sep 17 00:00:00 2001
From: Anshil Gandhi <95053726+gandhi56 at users.noreply.github.com>
Date: Thu, 16 Apr 2026 14:49:01 -0400
Subject: [PATCH 2/2] Update llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
Co-authored-by: Frederik Harwath <frederik at harwath.name>
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 30 +++++++++------------
1 file changed, 12 insertions(+), 18 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index 54d0dedd8e78d..32d043fa96a3e 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -3287,24 +3287,18 @@ void GCNPostGenericScheduler::schedNode(SUnit *SU, bool IsTopNode) {
// non-MFMA VALU immediately after an MFMA is modeled as costing an extra
// cycle (bumpCycle) so the post-scheduler's pressure/latency accounting
// matches hardware behavior for that transition.
- if (SU->isInstr()) {
- MachineInstr *MI = SU->getInstr();
- bool IsNonMFMAVALU = SIInstrInfo::isVALU(*MI) && !SIInstrInfo::isMFMA(*MI);
- if (IsNonMFMAVALU) {
- if (IsTopNode && LastTopScheduledIsMFMA) {
- Top.bumpCycle(Top.getCurrCycle() + 1);
- LastTopScheduledIsMFMA = false;
- }
- if (!IsTopNode && LastBottomScheduledIsMFMA) {
- Bot.bumpCycle(Bot.getCurrCycle() + 1);
- LastBottomScheduledIsMFMA = false;
- }
- } else if (SIInstrInfo::isMFMA(*MI)) {
- if (IsTopNode)
- LastTopScheduledIsMFMA = true;
- if (!IsTopNode)
- LastBottomScheduledIsMFMA = true;
- }
+ if (SU->isInstr() && SIInstrInfo::isVALU(*SU->getInstr())) {
+ bool IsMFMA = SIInstrInfo::isMFMA(*SU->getInstr());
+ auto Bump = [IsMFMA](SchedBoundary &Boundary, bool &LastIsMFMA) {
+ if (LastIsMFMA && !IsMFMA)
+ Boundary.bumpCycle(Boundary.getCurrCycle() + 1);
+ LastIsMFMA = IsMFMA;
+ };
+
+ if (IsTopNode)
+ Bump(Top, LastTopScheduledIsMFMA);
+ else
+ Bump(Bot, LastBottomScheduledIsMFMA);
}
PostGenericScheduler::schedNode(SU, IsTopNode);
More information about the llvm-commits
mailing list