[llvm] [AMDGPU] Add wait states between different MFMAs sharing an accumulator (PR #218363)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Aug 24 02:53:11 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Pankaj Dwivedi (PankajDwivedi-25)
<details>
<summary>Changes</summary>
A full-register src2/C read after an MFMA write emits no wait states,
relying on accumulator forwarding that only works while the chain stays
on one MFMA. Two different MFMAs sharing an accumulator instead need the
wait states of a partial overlap: on gfx950, v_mfma_f32_16x16x32_f16 then
v_mfma_f32_16x16x16_f16 on the same tuple needs 5 and got none. Compare
canonicalized opcodes so mac and register-bank forms still match.
Assisted-by: Claude Opus 5
---
Full diff: https://github.com/llvm/llvm-project/pull/218363.diff
3 Files Affected:
- (modified) llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp (+81-66)
- (modified) llvm/lib/Target/AMDGPU/GCNHazardRecognizer.h (+5)
- (modified) llvm/test/CodeGen/AMDGPU/mai-hazards-gfx942.mir (+30)
``````````diff
diff --git a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
index 9d01a529eb012..ea166d783a945 100644
--- a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
@@ -3164,6 +3164,16 @@ int GCNHazardRecognizer::checkMAIHazards908(MachineInstr *MI) const {
return WaitStatesNeeded;
}
+/// Map an MFMA opcode to a form that ignores the vdst/src2 register bank and
+/// the mac form, so two MFMAs can be compared for being the same instruction.
+static unsigned canonicalizeMFMAOpcode(unsigned Opc) {
+ if (int MacOp = AMDGPU::getMFMAEarlyClobberOp(Opc); MacOp != -1)
+ Opc = MacOp;
+ if (int AGPROp = AMDGPU::getAGPRFormOp(Opc); AGPROp != -1)
+ Opc = AGPROp;
+ return Opc;
+}
+
static int
GFX940_XDL_N_PassWritesVGPROverlappedXDLOrSMFMASrcCWaitStates(int NumPasses,
bool IsGFX950) {
@@ -3214,6 +3224,68 @@ static int GFX940_XDL_N_PassWritesVGPROverlappedSrcABWaitStates(int NumPasses,
return NumPasses + 3 + (NumPasses != 2 && IsGFX950);
}
+int GCNHazardRecognizer::getMFMAOverlappedSrcCWaitStates(
+ const MachineInstr *MI, const MachineInstr *MI1) const {
+ const int SMFMA4x4WritesVGPROverlappedSMFMASrcCWaitStates = 2;
+ const int SMFMA16x16WritesVGPROverlappedSMFMASrcCWaitStates = 8;
+ const int SMFMA32x32WritesVGPROverlappedSMFMASrcCWaitStates = 16;
+ const int SMFMA4x4WritesVGPROverlappedDMFMASrcCWaitStates = 3;
+ const int SMFMA16x16WritesVGPROverlappedDMFMASrcCWaitStates = 9;
+ const int SMFMA32x32WritesVGPROverlappedDMFMASrcCWaitStates = 17;
+ const int DMFMA16x16WritesVGPROverlappedSrcCWaitStates = 9;
+ const int GFX950_DMFMA16x16WritesVGPROverlappedSrcCWaitStates = 17;
+ const int DMFMA4x4WritesVGPROverlappedSrcCWaitStates = 4;
+
+ switch (MI1->getOpcode()) {
+ case AMDGPU::V_MFMA_F64_16X16X4F64_e64:
+ case AMDGPU::V_MFMA_F64_16X16X4F64_vgprcd_e64:
+ case AMDGPU::V_MFMA_F64_16X16X4F64_mac_e64:
+ case AMDGPU::V_MFMA_F64_16X16X4F64_mac_vgprcd_e64:
+ if (TII.isXDL(*MI))
+ return 0;
+ return ST.hasGFX950Insts()
+ ? GFX950_DMFMA16x16WritesVGPROverlappedSrcCWaitStates
+ : DMFMA16x16WritesVGPROverlappedSrcCWaitStates;
+ case AMDGPU::V_MFMA_F64_4X4X4F64_e64:
+ case AMDGPU::V_MFMA_F64_4X4X4F64_vgprcd_e64:
+ if (TII.isXDL(*MI))
+ return 0;
+ return DMFMA4x4WritesVGPROverlappedSrcCWaitStates;
+ default:
+ break;
+ }
+
+ int NumPasses = TSchedModel.computeInstrLatency(MI1);
+ if (ST.hasGFX940Insts()) {
+ if (TII.isXDL(*MI) && !TII.isXDL(*MI1))
+ return 0;
+
+ if (!TII.isXDL(*MI1))
+ return GFX940_SMFMA_N_PassWritesVGPROverlappedSMFMASrcCWaitStates(
+ NumPasses);
+ return TII.isXDL(*MI)
+ ? GFX940_XDL_N_PassWritesVGPROverlappedXDLOrSMFMASrcCWaitStates(
+ NumPasses, ST.hasGFX950Insts())
+ : GFX940_XDL_N_PassWritesVGPROverlappedSGEMMDGEMMSrcCWaitStates(
+ NumPasses, ST.hasGFX950Insts());
+ }
+
+ bool IsDGEMM = SIInstrInfo::isDGEMM(MI->getOpcode());
+ switch (NumPasses) {
+ case 2:
+ return IsDGEMM ? SMFMA4x4WritesVGPROverlappedDMFMASrcCWaitStates
+ : SMFMA4x4WritesVGPROverlappedSMFMASrcCWaitStates;
+ case 8:
+ return IsDGEMM ? SMFMA16x16WritesVGPROverlappedDMFMASrcCWaitStates
+ : SMFMA16x16WritesVGPROverlappedSMFMASrcCWaitStates;
+ case 16:
+ return IsDGEMM ? SMFMA32x32WritesVGPROverlappedDMFMASrcCWaitStates
+ : SMFMA32x32WritesVGPROverlappedSMFMASrcCWaitStates;
+ default:
+ llvm_unreachable("unexpected number of passes");
+ }
+}
+
int GCNHazardRecognizer::checkMAIHazards90A(MachineInstr *MI) const {
int WaitStatesNeeded = 0;
unsigned Opc = MI->getOpcode();
@@ -3242,15 +3314,6 @@ int GCNHazardRecognizer::checkMAIHazards90A(MachineInstr *MI) const {
// Loop for both DGEMM and S/HGEMM 2nd instruction.
for (const MachineOperand &Use : MI->explicit_uses()) {
const int LegacyVALUNotDotWritesVGPRWaitStates = 2;
- const int SMFMA4x4WritesVGPROverlappedSMFMASrcCWaitStates = 2;
- const int SMFMA16x16WritesVGPROverlappedSMFMASrcCWaitStates = 8;
- const int SMFMA32x32WritesVGPROverlappedSMFMASrcCWaitStates = 16;
- const int SMFMA4x4WritesVGPROverlappedDMFMASrcCWaitStates = 3;
- const int SMFMA16x16WritesVGPROverlappedDMFMASrcCWaitStates = 9;
- const int SMFMA32x32WritesVGPROverlappedDMFMASrcCWaitStates = 17;
- const int DMFMA16x16WritesVGPROverlappedSrcCWaitStates = 9;
- const int GFX950_DMFMA16x16WritesVGPROverlappedSrcCWaitStates = 17;
- const int DMFMA4x4WritesVGPROverlappedSrcCWaitStates = 4;
const int SMFMA4x4WritesVGPROverlappedSrcABWaitStates = 5;
const int SMFMA16x16WritesVGPROverlappedSrcABWaitStates = 11;
const int SMFMA32x32WritesVGPROverlappedSrcABWaitStates = 19;
@@ -3302,64 +3365,16 @@ int GCNHazardRecognizer::checkMAIHazards90A(MachineInstr *MI) const {
else if (ST.hasGFX940Insts() &&
TSchedModel.computeInstrLatency(MI1) == 2)
NeedWaitStates = GFX940_SMFMA4x4WritesVGPRFullSrcCWaitStates;
- } else {
- switch (Opc1) {
- case AMDGPU::V_MFMA_F64_16X16X4F64_e64:
- case AMDGPU::V_MFMA_F64_16X16X4F64_vgprcd_e64:
- case AMDGPU::V_MFMA_F64_16X16X4F64_mac_e64:
- case AMDGPU::V_MFMA_F64_16X16X4F64_mac_vgprcd_e64:
- if (!TII.isXDL(*MI))
- NeedWaitStates =
- ST.hasGFX950Insts()
- ? GFX950_DMFMA16x16WritesVGPROverlappedSrcCWaitStates
- : DMFMA16x16WritesVGPROverlappedSrcCWaitStates;
- break;
- case AMDGPU::V_MFMA_F64_4X4X4F64_e64:
- case AMDGPU::V_MFMA_F64_4X4X4F64_vgprcd_e64:
- if (!TII.isXDL(*MI))
- NeedWaitStates = DMFMA4x4WritesVGPROverlappedSrcCWaitStates;
- break;
- default:
- int NumPasses = TSchedModel.computeInstrLatency(MI1);
- if (ST.hasGFX940Insts()) {
- if (TII.isXDL(*MI) && !TII.isXDL(*MI1))
- break;
-
- NeedWaitStates =
- TII.isXDL(*MI1)
- ? (TII.isXDL(*MI)
- ? GFX940_XDL_N_PassWritesVGPROverlappedXDLOrSMFMASrcCWaitStates(
- NumPasses, ST.hasGFX950Insts())
- : GFX940_XDL_N_PassWritesVGPROverlappedSGEMMDGEMMSrcCWaitStates(
- NumPasses, ST.hasGFX950Insts()))
- : GFX940_SMFMA_N_PassWritesVGPROverlappedSMFMASrcCWaitStates(
- NumPasses);
- break;
- }
- switch (NumPasses) {
- case 2:
- NeedWaitStates =
- SIInstrInfo::isDGEMM(Opc)
- ? SMFMA4x4WritesVGPROverlappedDMFMASrcCWaitStates
- : SMFMA4x4WritesVGPROverlappedSMFMASrcCWaitStates;
- break;
- case 8:
- NeedWaitStates =
- SIInstrInfo::isDGEMM(Opc)
- ? SMFMA16x16WritesVGPROverlappedDMFMASrcCWaitStates
- : SMFMA16x16WritesVGPROverlappedSMFMASrcCWaitStates;
- break;
- case 16:
- NeedWaitStates =
- SIInstrInfo::isDGEMM(Opc)
- ? SMFMA32x32WritesVGPROverlappedDMFMASrcCWaitStates
- : SMFMA32x32WritesVGPROverlappedSMFMASrcCWaitStates;
- break;
- default:
- llvm_unreachable("unexpected number of passes");
- }
- }
+ // The accumulator forwarding path that allows zero wait states is only
+ // available while the chain stays on a single MFMA. Two different MFMAs
+ // sharing an accumulator need the wait states of a partial overlap.
+ if (ST.hasGFX940Insts() &&
+ canonicalizeMFMAOpcode(Opc) != canonicalizeMFMAOpcode(Opc1))
+ NeedWaitStates = std::max(NeedWaitStates,
+ getMFMAOverlappedSrcCWaitStates(MI, MI1));
+ } else {
+ NeedWaitStates = getMFMAOverlappedSrcCWaitStates(MI, MI1);
}
} else {
switch (Opc1) {
diff --git a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.h b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.h
index 422da557eb47d..8e30449b04fb3 100644
--- a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.h
+++ b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.h
@@ -242,6 +242,11 @@ class GCNHazardRecognizer final : public ScheduleHazardRecognizer {
int checkMAIHazards(MachineInstr *MI) const;
int checkMAIHazards908(MachineInstr *MI) const;
int checkMAIHazards90A(MachineInstr *MI) const;
+ /// Wait states needed when \p MI reads as src2/C a register that the earlier
+ /// MFMA \p MI1 wrote, taken from the wait state tables for overlapping
+ /// register tuples.
+ int getMFMAOverlappedSrcCWaitStates(const MachineInstr *MI,
+ const MachineInstr *MI1) const;
/// Pad the latency between neighboring MFMA instructions with s_nops. The
/// percentage of wait states to fill with s_nops is specified by the command
/// line option '-amdgpu-mfma-padding-ratio'.
diff --git a/llvm/test/CodeGen/AMDGPU/mai-hazards-gfx942.mir b/llvm/test/CodeGen/AMDGPU/mai-hazards-gfx942.mir
index 5a08e429283a8..99fd94ff772cb 100644
--- a/llvm/test/CodeGen/AMDGPU/mai-hazards-gfx942.mir
+++ b/llvm/test/CodeGen/AMDGPU/mai-hazards-gfx942.mir
@@ -2512,3 +2512,33 @@ body: |
$agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31 = V_MFMA_F32_32X32X4F16_e64 $vgpr126_vgpr127, $vgpr128_vgpr129, $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31, 0, 0, 0, implicit $mode, implicit $exec
$agpr2_agpr3_agpr4_agpr5 = V_SMFMAC_F32_16X16X32_F16_e64 $vgpr0_vgpr1, $vgpr2_vgpr3_vgpr4_vgpr5, $vgpr32, 0, 0, $agpr2_agpr3_agpr4_agpr5, implicit $mode, implicit $exec
...
+# GCN-LABEL: name: xdl_sgemm16x16_mfma_write_agpr_different_xdl_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_16X16X32_F16
+# GFX942-NEXT: S_NOP 4
+# GFX950-NEXT: S_NOP 5
+# GCN-NEXT: V_MFMA_F32_16X16X16F16
+name: xdl_sgemm16x16_mfma_write_agpr_different_xdl_mfma_read_same_agpr_as_srcc
+body: |
+ bb.0:
+ $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_16X16X32_F16_e64 $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5_vgpr6_vgpr7, $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+ $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_16X16X16F16_e64 $vgpr0_vgpr1, $vgpr2_vgpr3, $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+...
+# GCN-LABEL: name: nonxdl_sgemm16x16_mfma_write_agpr_different_nonxdl_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_16X16X4F32
+# GCN-NEXT: S_NOP 7
+# GCN-NEXT: V_MFMA_F32_4X4X1F32
+name: nonxdl_sgemm16x16_mfma_write_agpr_different_nonxdl_mfma_read_same_agpr_as_srcc
+body: |
+ bb.0:
+ $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_16X16X4F32_e64 $vgpr0, $vgpr1, $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+ $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X1F32_e64 $vgpr0, $vgpr1, $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+...
+# GCN-LABEL: name: xdl_sgemm16x16_mac_mfma_write_agpr_nonmac_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_16X16X4F16_mac
+# GCN-NEXT: V_MFMA_F32_16X16X4F16
+name: xdl_sgemm16x16_mac_mfma_write_agpr_nonmac_mfma_read_same_agpr_as_srcc
+body: |
+ bb.0:
+ $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15 = V_MFMA_F32_16X16X4F16_mac_e64 $vgpr0_vgpr1, $vgpr2_vgpr3, $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15, 0, 0, 0, implicit $mode, implicit $exec
+ $agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31 = V_MFMA_F32_16X16X4F16_e64 $vgpr0_vgpr1, $vgpr2_vgpr3, $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15, 0, 0, 0, implicit $mode, implicit $exec
+...
``````````
</details>
https://github.com/llvm/llvm-project/pull/218363
More information about the llvm-commits
mailing list