[llvm] 4362da2 - [AMDGPU] Add wait states between different MFMAs sharing an accumulator (#218363)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Sep 28 00:40:22 PDT 2026
Author: Pankaj Dwivedi
Date: 2026-09-28T13:10:15+05:30
New Revision: 4362da25049882128f50a0f6b4fb69c1a0f8fd57
URL: https://github.com/llvm/llvm-project/commit/4362da25049882128f50a0f6b4fb69c1a0f8fd57
DIFF: https://github.com/llvm/llvm-project/commit/4362da25049882128f50a0f6b4fb69c1a0f8fd57.diff
LOG: [AMDGPU] Add wait states between different MFMAs sharing an accumulator (#218363)
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
Added:
Modified:
llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
llvm/lib/Target/AMDGPU/GCNHazardRecognizer.h
llvm/test/CodeGen/AMDGPU/mai-hazards-gfx90a.mir
llvm/test/CodeGen/AMDGPU/mai-hazards-gfx942.mir
Removed:
################################################################################
diff --git a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
index 08c3d7dd9b098..d9e52f697ba66 100644
--- a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
@@ -3158,6 +3158,24 @@ int GCNHazardRecognizer::checkMAIHazards908(MachineInstr *MI) const {
return WaitStatesNeeded;
}
+/// One MFMA can be written with up to four opcodes that
diff er only in how vdst
+/// and src2 are encoded: both are either AGPRs or VGPRs, and the mac form ties
+/// vdst to src2 instead of taking them as separate operands. \returns the AGPR,
+/// non-mac opcode, so that every form of the same MFMA maps to one value.
+static unsigned getMFMANonMacAGPRFormOp(unsigned Opc) {
+ if (int NonMacOp = AMDGPU::getMFMAEarlyClobberOp(Opc); NonMacOp != -1)
+ Opc = NonMacOp;
+ if (int AGPROp = AMDGPU::getAGPRFormOp(Opc); AGPROp != -1)
+ Opc = AGPROp;
+ return Opc;
+}
+
+/// \returns true if \p Opc0 and \p Opc1 are the same MFMA, ignoring the mac
+/// form and whether vdst/src2 are AGPRs or VGPRs.
+static bool isSameMFMA(unsigned Opc0, unsigned Opc1) {
+ return getMFMANonMacAGPRFormOp(Opc0) == getMFMANonMacAGPRFormOp(Opc1);
+}
+
static int
GFX940_XDL_N_PassWritesVGPROverlappedXDLOrSMFMASrcCWaitStates(int NumPasses,
bool IsGFX950) {
@@ -3208,6 +3226,66 @@ static int GFX940_XDL_N_PassWritesVGPROverlappedSrcABWaitStates(int NumPasses,
return NumPasses + 3 + (NumPasses != 2 && IsGFX950);
}
+int GCNHazardRecognizer::getMFMAOverlappedSrcCWaitStates(
+ const MachineInstr *Reader, const MachineInstr *Writer) const {
+ constexpr int SMFMA4x4WritesVGPROverlappedSMFMASrcCWaitStates = 2;
+ constexpr int SMFMA16x16WritesVGPROverlappedSMFMASrcCWaitStates = 8;
+ constexpr int SMFMA32x32WritesVGPROverlappedSMFMASrcCWaitStates = 16;
+ constexpr int SMFMA4x4WritesVGPROverlappedDMFMASrcCWaitStates = 3;
+ constexpr int SMFMA16x16WritesVGPROverlappedDMFMASrcCWaitStates = 9;
+ constexpr int SMFMA32x32WritesVGPROverlappedDMFMASrcCWaitStates = 17;
+ constexpr int DMFMA16x16WritesVGPROverlappedSrcCWaitStates = 9;
+ constexpr int GFX950_DMFMA16x16WritesVGPROverlappedSrcCWaitStates = 17;
+ constexpr int DMFMA4x4WritesVGPROverlappedSrcCWaitStates = 4;
+
+ // An XDL read of a non-XDL result needs no wait states. DGEMM is never XDL,
+ // so this also covers the f64 writers handled below.
+ if (TII.isXDL(*Reader) && !TII.isXDL(*Writer))
+ return 0;
+
+ switch (Writer->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:
+ return ST.hasGFX950Insts()
+ ? GFX950_DMFMA16x16WritesVGPROverlappedSrcCWaitStates
+ : DMFMA16x16WritesVGPROverlappedSrcCWaitStates;
+ case AMDGPU::V_MFMA_F64_4X4X4F64_e64:
+ case AMDGPU::V_MFMA_F64_4X4X4F64_vgprcd_e64:
+ return DMFMA4x4WritesVGPROverlappedSrcCWaitStates;
+ default:
+ break;
+ }
+
+ int NumPasses = TSchedModel.computeInstrLatency(Writer);
+ if (ST.hasGFX940Insts()) {
+ if (!TII.isXDL(*Writer))
+ return GFX940_SMFMA_N_PassWritesVGPROverlappedSMFMASrcCWaitStates(
+ NumPasses);
+ return TII.isXDL(*Reader)
+ ? GFX940_XDL_N_PassWritesVGPROverlappedXDLOrSMFMASrcCWaitStates(
+ NumPasses, ST.hasGFX950Insts())
+ : GFX940_XDL_N_PassWritesVGPROverlappedSGEMMDGEMMSrcCWaitStates(
+ NumPasses, ST.hasGFX950Insts());
+ }
+
+ bool IsDGEMM = SIInstrInfo::isDGEMM(Reader->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();
@@ -3236,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;
@@ -3298,64 +3367,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
diff erent MFMAs
+ // sharing an accumulator need the wait states of a partial overlap.
+ if (ST.hasGFX940Insts() && !isSameMFMA(Opc, 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..705c372f5f007 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 MFMA \p Reader reads as src2/C a register that the
+ /// earlier MFMA \p Writer wrote, taken from the wait state tables for
+ /// overlapping register tuples.
+ int getMFMAOverlappedSrcCWaitStates(const MachineInstr *Reader,
+ const MachineInstr *Writer) 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-gfx90a.mir b/llvm/test/CodeGen/AMDGPU/mai-hazards-gfx90a.mir
index 5fb6824d41dd0..bf11056022b7d 100644
--- a/llvm/test/CodeGen/AMDGPU/mai-hazards-gfx90a.mir
+++ b/llvm/test/CodeGen/AMDGPU/mai-hazards-gfx90a.mir
@@ -1444,3 +1444,12 @@ body: |
$agpr0_agpr1 = V_MFMA_F64_4X4X4F64_e64 $agpr0_agpr1, $agpr0_agpr1, $agpr0_agpr1, 0, 0, 0, implicit $mode, implicit $exec
BUFFER_STORE_DWORDX2_OFFEN_exact $vgpr2_vgpr3, $vgpr0, $sgpr0_sgpr1_sgpr2_sgpr3, 0, 0, 0, 0, implicit $exec
...
+# GCN-LABEL: name: xdl_8pass_mfma_write_agpr_
diff erent_xdl_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_16X16X4F16
+# GCN-NEXT: V_MFMA_I32_16X16X4I8
+name: xdl_8pass_mfma_write_agpr_
diff erent_xdl_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_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
+ $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15 = V_MFMA_I32_16X16X4I8_e64 $vgpr0, $vgpr1, $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15, 0, 0, 0, implicit $mode, implicit $exec
+...
diff --git a/llvm/test/CodeGen/AMDGPU/mai-hazards-gfx942.mir b/llvm/test/CodeGen/AMDGPU/mai-hazards-gfx942.mir
index bdc96860632a4..2f78efbdd256e 100644
--- a/llvm/test/CodeGen/AMDGPU/mai-hazards-gfx942.mir
+++ b/llvm/test/CodeGen/AMDGPU/mai-hazards-gfx942.mir
@@ -2557,3 +2557,161 @@ body: |
S_NOP 2
$vgpr0 = V_MOV_B32_e32 1, implicit $exec
...
+# GCN-LABEL: name: xdl_sgemm16x16_mfma_write_agpr_
diff erent_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_
diff erent_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_
diff erent_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_
diff erent_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
+...
+# GCN-LABEL: name: xdl_2pass_mfma_write_agpr_
diff erent_xdl_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_4X4X4F16
+# GFX942-NEXT: S_NOP 2
+# GFX950-NEXT: S_NOP 3
+# GCN-NEXT: V_MFMA_F32_16X16X16F16
+name: xdl_2pass_mfma_write_agpr_
diff erent_xdl_mfma_read_same_agpr_as_srcc
+body: |
+ bb.0:
+ $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X4F16_e64 $vgpr0_vgpr1, $vgpr2_vgpr3, $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: xdl_8pass_mfma_write_agpr_
diff erent_xdl_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_16X16X4F16
+# GFX942-NEXT: S_NOP 8
+# GFX950-NEXT: S_NOP 9
+# GCN-NEXT: V_MFMA_I32_16X16X4I8
+name: xdl_8pass_mfma_write_agpr_
diff erent_xdl_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_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
+ $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15 = V_MFMA_I32_16X16X4I8_e64 $vgpr0, $vgpr1, $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15, 0, 0, 0, implicit $mode, implicit $exec
+...
+# GCN-LABEL: name: xdl_16pass_mfma_write_agpr_nonxdl_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_32X32X4F16
+# GCN-NEXT: S_NOP 15
+# GFX942-NEXT: S_NOP 0
+# GFX950-NEXT: S_NOP 1
+# GCN-NEXT: V_MFMA_F32_32X32X1F32
+name: xdl_16pass_mfma_write_agpr_nonxdl_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_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31 = V_MFMA_F32_32X32X4F16_e64 $vgpr0_vgpr1, $vgpr2_vgpr3, $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
+ $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_32X32X1F32_e64 $vgpr0, $vgpr1, $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
+...
+# GCN-LABEL: name: xdl_8pass_mfma_write_agpr_nonxdl_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_16X16X4F16
+# GFX942-NEXT: S_NOP 8
+# GFX950-NEXT: S_NOP 9
+# GCN-NEXT: V_MFMA_F32_16X16X1F32
+name: xdl_8pass_mfma_write_agpr_nonxdl_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_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
+ $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15 = V_MFMA_F32_16X16X1F32_e64 $vgpr0, $vgpr1, $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15, 0, 0, 0, implicit $mode, implicit $exec
+...
+# GCN-LABEL: name: xdl_2pass_mfma_write_agpr_nonxdl_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_4X4X4F16
+# GCN-NEXT: S_NOP 2
+# GCN-NEXT: V_MFMA_F32_16X16X4F32
+name: xdl_2pass_mfma_write_agpr_nonxdl_mfma_read_same_agpr_as_srcc
+body: |
+ bb.0:
+ $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X4F16_e64 $vgpr0_vgpr1, $vgpr2_vgpr3, $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+ $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_16X16X4F32_e64 $vgpr0, $vgpr1, $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+...
+# GCN-LABEL: name: nonxdl_8pass_mfma_write_agpr_xdl_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_16X16X4F32
+# GCN-NEXT: V_MFMA_F32_16X16X16F16
+name: nonxdl_8pass_mfma_write_agpr_xdl_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_16X16X16F16_e64 $vgpr0_vgpr1, $vgpr2_vgpr3, $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+...
+# GCN-LABEL: name: nonxdl_2pass_mfma_write_agpr_xdl_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_4X4X1F32
+# GCN-NEXT: S_NOP 1
+# GCN-NEXT: V_MFMA_F32_16X16X16F16
+name: nonxdl_2pass_mfma_write_agpr_xdl_mfma_read_same_agpr_as_srcc
+body: |
+ bb.0:
+ $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X1F32_e64 $vgpr0, $vgpr1, $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: xdl_4pass_mfma_write_agpr_smfmac_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_16X16X16F16
+# GFX942-NEXT: S_NOP 4
+# GFX950-NEXT: S_NOP 5
+# GCN-NEXT: V_SMFMAC_F32_16X16X32_F16
+name: xdl_4pass_mfma_write_agpr_smfmac_read_same_agpr_as_srcc
+body: |
+ bb.0:
+ $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_16X16X16F16_e64 $vgpr0_vgpr1, $vgpr2_vgpr3, $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+ $agpr0_agpr1_agpr2_agpr3 = V_SMFMAC_F32_16X16X32_F16_e64 $vgpr0_vgpr1, $vgpr2_vgpr3_vgpr4_vgpr5, $vgpr32, 0, 0, $agpr0_agpr1_agpr2_agpr3, implicit $mode, implicit $exec
+...
+# GCN-LABEL: name: xdl_16pass_mfma_write_agpr_
diff erent_xdl_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_32X32X4F16
+# GCN-NEXT: S_NOP 15
+# GFX942-NEXT: S_NOP 0
+# GFX950-NEXT: S_NOP 1
+# GCN-NEXT: V_MFMA_F32_32X32X4BF16_1K
+name: xdl_16pass_mfma_write_agpr_
diff erent_xdl_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_agpr16_agpr17_agpr18_agpr19_agpr20_agpr21_agpr22_agpr23_agpr24_agpr25_agpr26_agpr27_agpr28_agpr29_agpr30_agpr31 = V_MFMA_F32_32X32X4F16_e64 $vgpr0_vgpr1, $vgpr2_vgpr3, $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
+ $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_32X32X4BF16_1K_e64 $vgpr0_vgpr1, $vgpr2_vgpr3, $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
+...
+# GCN-LABEL: name: xdl_4pass_mfma_write_agpr_nonxdl_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_16X16X16F16
+# GFX942-NEXT: S_NOP 4
+# GFX950-NEXT: S_NOP 5
+# GCN-NEXT: V_MFMA_F32_16X16X4F32
+name: xdl_4pass_mfma_write_agpr_nonxdl_mfma_read_same_agpr_as_srcc
+body: |
+ bb.0:
+ $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_16X16X16F16_e64 $vgpr0_vgpr1, $vgpr2_vgpr3, $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+ $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_16X16X4F32_e64 $vgpr0, $vgpr1, $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+...
+# GCN-LABEL: name: nonxdl_2pass_mfma_write_agpr_
diff erent_nonxdl_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_4X4X1F32
+# GCN-NEXT: S_NOP 1
+# GCN-NEXT: V_MFMA_F32_16X16X4F32
+name: nonxdl_2pass_mfma_write_agpr_
diff erent_nonxdl_mfma_read_same_agpr_as_srcc
+body: |
+ bb.0:
+ $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X1F32_e64 $vgpr0, $vgpr1, $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+ $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_16X16X4F32_e64 $vgpr0, $vgpr1, $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+...
+# GCN-LABEL: name: nonxdl_16pass_mfma_write_agpr_
diff erent_nonxdl_mfma_read_same_agpr_as_srcc
+# GCN: V_MFMA_F32_32X32X2F32
+# GCN-NEXT: S_NOP 15
+# GCN-NEXT: V_MFMA_F32_16X16X1F32
+name: nonxdl_16pass_mfma_write_agpr_
diff erent_nonxdl_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_32X32X2F32_e64 $vgpr0, $vgpr1, $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15, 0, 0, 0, implicit $mode, implicit $exec
+ $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15 = V_MFMA_F32_16X16X1F32_e64 $vgpr0, $vgpr1, $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15, 0, 0, 0, implicit $mode, implicit $exec
+...
More information about the llvm-commits
mailing list