[llvm] [AMDGPU] Add initial version of gfx950 MFMA co-exec rules (PR #214682)
Romanov Vlad via llvm-commits
llvm-commits at lists.llvm.org
Tue Aug 25 02:39:41 PDT 2026
https://github.com/romanovvlad updated https://github.com/llvm/llvm-project/pull/214682
>From 4e83b53d66589df40f39f507831418615f990455 Mon Sep 17 00:00:00 2001
From: Vlad <Vladislav.Romanov at amd.com>
Date: Fri, 7 Aug 2026 03:40:52 -0500
Subject: [PATCH 01/11] [AMDGPU] Add initial version of gfx950 MFMA co-exec
rules
Added co-exec rules for all MFMA instructions which were added in
`gfx950`. Legacy instructions which were inherited from older Instinct
GPUs generations use fallback for now.
Added rules aren't exactly precise and serve merely as a baseline for now.
Also remove the warning when running coexec sched on gfx950.
Co-authored-by: Alexey Sachkov <alexey.sachkov at amd.com>
---
llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h | 170 +++++++++++++++++-
.../AMDGPU/AMDGPUCoExecSchedStrategy.cpp | 11 +-
.../lib/Target/AMDGPU/AMDGPUTargetMachine.cpp | 5 +-
.../lib/Target/AMDGPU/GCNHazardRecognizer.cpp | 20 ++-
.../AMDGPU/coexec-mfma-interleave-gfx950.mir | 48 +++++
.../CodeGen/AMDGPU/coexec-rewrite-mfma.ll | 2 -
6 files changed, 244 insertions(+), 12 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.mir
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
index 6f0f50ed7472b..a07eb8ea74427 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
@@ -50,7 +50,7 @@ enum class CoExecMask : uint16_t {
DS = 1 << 4, // LDS read/write
VMEM = 1 << 5, // Global memory
SMEM = 1 << 6, // Scalar memory
- WMMA = 1 << 7, // Next WMMA (V stages only)
+ WMMA = 1 << 7, // Next WMMA (V stages only), or MFMA
All = 0xFFFF,
MEM = DS | VMEM | SMEM,
@@ -346,12 +346,180 @@ inline CoExecInfo CoExecInfo::build(unsigned TotalWindow, const char *Pattern) {
return Info;
}
+/// Get co-execution info for a gfx950 MFMA instruction.
+/// The occupancy (cycles until the next MFMA may issue) is expressed as the
+/// first stage carrying the WMMA bit.
+inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
+ CoExecInfo Res;
+ for (unsigned I = 0; I < MaxCoExecStages; ++I)
+ Res.Slots[I].Mask = CoExecMask::None;
+
+ // TODO: Implement proper patterns support (for debugging purposes).
+ // Existing pattern letters are WMMA-specific and will probably be confusing
+ // if used as-is for MFMA. Inventing new MFMA-specific letters is an option,
+ // but perhaps the pattern should be instead dynamically reconstructed when
+ // needed by printing specific slots in full instead of a key for them.
+ Res.Pattern = "undefinedundefinedundefinedundefined";
+
+ // MFMA co-exec slots are incremental, i.e. for every slot N it supports all
+ // instructions which were supported by the previous slot N-1 and may support
+ // something extra.
+ auto AllowCoExec = [](CoExecInfo &Info, CoExecMaskT ExtraBits,
+ unsigned StartIndex) {
+ for (unsigned Index = StartIndex; Index < Info.TotalWindow; ++Index)
+ Info.Slots[Index].Mask |= ExtraBits;
+ };
+
+ switch (MI.getOpcode()) {
+ // 4-cycle occupancy, 8-cycle window.
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f6_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f6_vgprcd_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f4_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f4_vgprcd_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f6_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f6_vgprcd_e64:
+ case V_MFMA_F32_16X16X32_BF16_e64:
+ case V_MFMA_F32_16X16X32_BF16_vgprcd_e64:
+ case V_MFMA_I32_16X16X64_I8_e64:
+ case V_MFMA_I32_16X16X64_I8_vgprcd_e64:
+ case V_MFMA_F32_16X16X32_F16_e64:
+ case V_MFMA_F32_16X16X32_F16_vgprcd_e64:
+ // Distinction marker to simplify mapping to the programming guide
+ case V_SMFMAC_F32_16X16X64_BF16_e64:
+ case V_SMFMAC_I32_16X16X128_I8_e64:
+ case V_SMFMAC_F32_16X16X128_BF8_BF8_e64:
+ case V_SMFMAC_F32_16X16X128_BF8_FP8_e64:
+ case V_SMFMAC_F32_16X16X128_FP8_BF8_e64:
+ case V_SMFMAC_F32_16X16X128_FP8_FP8_e64:
+ case V_SMFMAC_F32_16X16X64_F16_e64:
+ Res.TotalWindow = 8;
+ AllowCoExec(Res, CoExecMask::SALU, 1);
+ AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU, 2);
+ AllowCoExec(Res, CoExecMask::WMMA, 4);
+ return Res;
+
+ // 8-cycle occupancy, 12-cycle window.
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f8_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f8_vgprcd_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f8_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f8_vgprcd_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f4_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f4_vgprcd_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_vgprcd_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64:
+ Res.TotalWindow = 12;
+ AllowCoExec(Res, CoExecMask::SALU, 1);
+ AllowCoExec(Res, CoExecMask::DS, 2);
+ AllowCoExec(Res, CoExecMask::VALU, 3);
+ AllowCoExec(Res, CoExecMask::WMMA, 8);
+ return Res;
+
+ // 4-cycle occupancy, 8-cycle window.
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_vgprcd_e64:
+ case V_MFMA_F32_32X32X16_BF16_e64:
+ case V_MFMA_F32_32X32X16_BF16_mac_e64:
+ case V_MFMA_F32_32X32X16_BF16_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X16_BF16_vgprcd_e64:
+ case V_MFMA_I32_32X32X16I8_e64:
+ case V_MFMA_I32_32X32X16I8_mac_e64:
+ case V_MFMA_I32_32X32X16I8_mac_vgprcd_e64:
+ case V_MFMA_I32_32X32X16I8_vgprcd_e64:
+ case V_MFMA_F32_32X32X16_F16_e64:
+ case V_MFMA_F32_32X32X16_F16_mac_e64:
+ case V_MFMA_F32_32X32X16_F16_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X16_F16_vgprcd_e64:
+ Res.TotalWindow = 8;
+ AllowCoExec(Res, CoExecMask::SALU, 1);
+ AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU, 2);
+ AllowCoExec(Res, CoExecMask::WMMA, 4);
+ return Res;
+
+ // 16-cycle occupancy, 20-cycle window.
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_vgprcd_e64:
+ Res.TotalWindow = 20;
+ AllowCoExec(Res, CoExecMask::SALU, 1);
+ AllowCoExec(Res, CoExecMask::DS, 2);
+ AllowCoExec(Res, CoExecMask::VALU, 3);
+ AllowCoExec(Res, CoExecMask::WMMA, 16);
+ return Res;
+
+ // 9-cycle occupancy, 12-cycle window.
+ case V_SMFMAC_F32_32X32X32_BF16_e64:
+ case V_SMFMAC_I32_32X32X64_I8_e64:
+ case V_SMFMAC_F32_32X32X64_BF8_BF8_e64:
+ case V_SMFMAC_F32_32X32X64_BF8_FP8_e64:
+ case V_SMFMAC_F32_32X32X64_FP8_BF8_e64:
+ case V_SMFMAC_F32_32X32X64_FP8_FP8_e64:
+ case V_SMFMAC_F32_32X32X32_F16_e64:
+ Res.TotalWindow = 12;
+ AllowCoExec(Res, CoExecMask::SALU, 1);
+ AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU, 4);
+ AllowCoExec(Res, CoExecMask::WMMA, 9);
+ return Res;
+
+ // 18-cycle occupancy, 19-cycle window.
+ case V_MFMA_F64_16X16X4F64_e64:
+ case V_MFMA_F64_16X16X4F64_mac_e64:
+ case V_MFMA_F64_16X16X4F64_mac_vgprcd_e64:
+ case V_MFMA_F64_16X16X4F64_vgprcd_e64:
+ Res.TotalWindow = 19;
+ AllowCoExec(Res, CoExecMask::DS | CoExecMask::SALU, 0);
+ AllowCoExec(Res, CoExecMask::WMMA | CoExecMask::VALU, 18);
+ return Res;
+
+ default:
+ // Default fallback: permissive 8-cycle pattern
+ return CoExecInfo::build(9, "AAAAAAAAA");
+ }
+}
+
/// Get co-execution info for a WMMA instruction, selecting the per-cycle slot
/// pattern from the opcode (and operand formats for the F8F6F4 variants).
inline CoExecInfo getCoExecInfo(const MachineInstr &MI,
const SIInstrInfo &TII) {
unsigned Opc = MI.getOpcode();
+ if (TII.isMFMA(MI))
+ return getMFMACoExecInfo(MI);
+
// Scaled variants (LD_SCALE rule) absorb the next WMMA in the last I slot.
bool HasScaling = AMDGPU::getHasMatrixScale(Opc);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
index 726be8c7f0982..c68e1b8508148 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
@@ -60,7 +60,16 @@ InstructionFlavor llvm::AMDGPU::classifyFlavor(const MachineInstr &MI,
if (SII.isLDSDMA(MI))
return InstructionFlavor::DMA;
- if (SII.isMFMAorWMMA(MI))
+ if (SII.isMFMA(MI)) {
+ // TODO: Consider further sub-classifying this (XDL, XDL2x, S/DGEMM).
+ // GFX9 SPG sub-classifies MFMA into XDL, XDL2x and S/DGEMM, because only
+ // certain sub-classes can be co-executed in certain slots. For now, we
+ // simply treat them all as one to simplify the change and leave the rest
+ // to a follow-up fine-tuning.
+ return InstructionFlavor::WMMA;
+ }
+
+ if (SII.isWMMA(MI) || SII.isSWMMAC(MI))
return InstructionFlavor::WMMA;
if (SII.isTRANS(MI))
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
index 30ea659cbf322..316e04c63f373 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -613,11 +613,12 @@ StringRef llvm::AMDGPU::getSchedStrategy(const Function &F) {
static void
diagnoseUnsupportedCoExecSchedulerSelection(const Function &F,
const GCNSubtarget &ST) {
- if (ST.hasGFX1250Insts())
+ if (ST.hasGFX1250Insts() || ST.hasGFX950Insts())
return;
F.getContext().diagnose(DiagnosticInfoUnsupported(
- F, "'amdgpu-sched-strategy'='coexec' is only supported for gfx1250",
+ F,
+ "'amdgpu-sched-strategy'='coexec' is only supported for gfx1250/gfx950",
DiagnosticLocation(), DS_Warning));
}
diff --git a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
index a3842f43a62fa..4d1b8a7e3a830 100644
--- a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
@@ -11,6 +11,7 @@
//===----------------------------------------------------------------------===//
#include "GCNHazardRecognizer.h"
+#include "AMDGPUTargetMachine.h"
#include "AMDGPUWaitcntUtils.h"
#include "GCNSubtarget.h"
#include "MCTargetDesc/AMDGPUMCTargetDesc.h"
@@ -216,16 +217,23 @@ bool GCNHazardRecognizer::hasCoExecWindowModel() const {
// gfx1251 and gfx12.5-generic report the same co-execution hazard features
// but have different WMMA latencies, so they need their own slot patterns
// before they can be modeled here.
- return ST.hasWMMACoexecutionHazards() && ST.hasTransCoexecutionHazard() &&
- AMDGPU::isGFX1250(ST);
+ if (ST.hasWMMACoexecutionHazards() && ST.hasTransCoexecutionHazard() &&
+ AMDGPU::isGFX1250(ST))
+ return true;
+
+ if (ST.hasGFX950Insts() &&
+ AMDGPU::getSchedStrategy(MF.getFunction()) == "coexec")
+ return true;
+
+ return false;
}
void GCNHazardRecognizer::updateWMMAWindowState(const MachineInstr &MI) {
if (!hasCoExecWindowModel())
return;
- // Check if this is a WMMA instruction.
- if (!SIInstrInfo::isWMMA(MI) && !SIInstrInfo::isSWMMAC(MI))
+ if (!SIInstrInfo::isWMMA(MI) && !SIInstrInfo::isSWMMAC(MI) &&
+ !SIInstrInfo::isMFMA(MI))
return;
// If a previous window was still active, dump it before starting a new one.
@@ -276,9 +284,9 @@ void GCNHazardRecognizer::updateMultiCycleVALUState(const MachineInstr &MI) {
if (!SIInstrInfo::isVALU(MI, /*AllowLDSDMA=*/true))
return;
- // Skip WMMA and TRANS - they have their own tracking.
+ // Skip WMMA, MFMA, and TRANS - they have their own tracking.
if (SIInstrInfo::isWMMA(MI) || SIInstrInfo::isSWMMAC(MI) ||
- SIInstrInfo::isTRANS(MI))
+ SIInstrInfo::isMFMA(MI) || SIInstrInfo::isTRANS(MI))
return;
unsigned RepeatRate = TII.getRepeatRate(MI);
diff --git a/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.mir b/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.mir
new file mode 100644
index 0000000000000..9a329f75132da
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.mir
@@ -0,0 +1,48 @@
+# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 6
+# RUN: llc -mtriple=amdgpu9.50-amd-amdhsa -run-pass=machine-scheduler -amdgpu-sched-strategy=coexec -verify-misched %s -o - | FileCheck %s
+
+# Test MFMA coexec window interleaving on gfx950.
+# Expects 2 VALUs + 1 SALU interleaved per MFMA shadow.
+
+---
+name: mfma_16x16_interleave
+tracksRegLiveness: true
+machineFunctionInfo:
+ isEntryFunction: true
+body: |
+ bb.0:
+
+ ; Three independent MFMAs
+ ; CHECK-LABEL: name: mfma_16x16_interleave
+ ; CHECK: %m0:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, undef %acc0:areg_128_align2, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %sa0:sgpr_32 = S_ADD_I32 undef $sgpr10, undef $sgpr11, implicit-def $scc
+ ; CHECK-NEXT: %a0:vgpr_32 = V_ADD_F32_e32 undef %v0:vgpr_32, undef %v1:vgpr_32, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %a1:vgpr_32 = V_ADD_F32_e32 undef %v1:vgpr_32, undef %v2:vgpr_32, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %m1:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, undef %acc1:areg_128_align2, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %sa1:sgpr_32 = S_ADD_I32 undef $sgpr12, undef $sgpr13, implicit-def $scc
+ ; CHECK-NEXT: %a2:vgpr_32 = V_ADD_F32_e32 undef %v2:vgpr_32, undef %v3:vgpr_32, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %a3:vgpr_32 = V_ADD_F32_e32 undef %v3:vgpr_32, undef %v4:vgpr_32, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %m2:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s2:vreg_128_align2, undef %s3:vreg_128_align2, undef %acc2:areg_128_align2, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %sa2:sgpr_32 = S_ADD_I32 undef $sgpr14, undef $sgpr15, implicit-def $scc
+ ; CHECK-NEXT: %a4:vgpr_32 = V_ADD_F32_e32 undef %v4:vgpr_32, undef %v5:vgpr_32, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %a5:vgpr_32 = V_ADD_F32_e32 undef %v5:vgpr_32, undef %v0:vgpr_32, implicit $mode, implicit $exec
+ ; CHECK-NEXT: S_ENDPGM 0, implicit %m0, implicit %m1, implicit %m2, implicit %a0, implicit %a1, implicit %a2, implicit %a3, implicit %a4, implicit %a5, implicit %sa0, implicit %sa1, implicit %sa2
+ %m0:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, undef %acc0:areg_128_align2, 0, 0, 0, implicit $mode, implicit $exec
+ %m1:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, undef %acc1:areg_128_align2, 0, 0, 0, implicit $mode, implicit $exec
+ %m2:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s2:vreg_128_align2, undef %s3:vreg_128_align2, undef %acc2:areg_128_align2, 0, 0, 0, implicit $mode, implicit $exec
+
+ ; Independent VALUs
+ %a0:vgpr_32 = V_ADD_F32_e32 undef %v0:vgpr_32, undef %v1:vgpr_32, implicit $mode, implicit $exec
+ %a1:vgpr_32 = V_ADD_F32_e32 undef %v1:vgpr_32, undef %v2:vgpr_32, implicit $mode, implicit $exec
+ %a2:vgpr_32 = V_ADD_F32_e32 undef %v2:vgpr_32, undef %v3:vgpr_32, implicit $mode, implicit $exec
+ %a3:vgpr_32 = V_ADD_F32_e32 undef %v3:vgpr_32, undef %v4:vgpr_32, implicit $mode, implicit $exec
+ %a4:vgpr_32 = V_ADD_F32_e32 undef %v4:vgpr_32, undef %v5:vgpr_32, implicit $mode, implicit $exec
+ %a5:vgpr_32 = V_ADD_F32_e32 undef %v5:vgpr_32, undef %v0:vgpr_32, implicit $mode, implicit $exec
+
+ ; Independent SALUs
+ %sa0:sgpr_32 = S_ADD_I32 undef $sgpr10, undef $sgpr11, implicit-def $scc
+ %sa1:sgpr_32 = S_ADD_I32 undef $sgpr12, undef $sgpr13, implicit-def $scc
+ %sa2:sgpr_32 = S_ADD_I32 undef $sgpr14, undef $sgpr15, implicit-def $scc
+
+ S_ENDPGM 0, implicit %m0, implicit %m1, implicit %m2, implicit %a0, implicit %a1, implicit %a2, implicit %a3, implicit %a4, implicit %a5, implicit %sa0, implicit %sa1, implicit %sa2
+...
diff --git a/llvm/test/CodeGen/AMDGPU/coexec-rewrite-mfma.ll b/llvm/test/CodeGen/AMDGPU/coexec-rewrite-mfma.ll
index 17e1012c05f6a..a632f5b3d161d 100644
--- a/llvm/test/CodeGen/AMDGPU/coexec-rewrite-mfma.ll
+++ b/llvm/test/CodeGen/AMDGPU/coexec-rewrite-mfma.ll
@@ -7,8 +7,6 @@
; exact) MFMAs in a phi-carried loop with high VGPR pressure, every MFMA in the
; loop is rewritten to the AGPR-destination form (v_mfma_f32_16x16x32_f16 a*).
-; FIXME: To enable on gfx950?
-; CHECK: warning: {{.*}}'amdgpu-sched-strategy'='coexec' is only supported for gfx1250
; CHECK-LABEL: v5_local_prefetch:
; CHECK-COUNT-28: v_mfma_f32_16x16x32_f16 a
; CHECK-NOT: v_mfma_f32_16x16x32_f16 v
>From 80667305018ee76d587c9d9f28e2b22cdad9d25c Mon Sep 17 00:00:00 2001
From: Alexey Sachkov <alexey.sachkov at amd.com>
Date: Tue, 18 Aug 2026 10:15:06 -0500
Subject: [PATCH 02/11] Clarify comment
---
llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h | 4 +++-
1 file changed, 3 insertions(+), 1 deletion(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
index a07eb8ea74427..7a7d424d0d3f0 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
@@ -386,7 +386,9 @@ inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
case V_MFMA_I32_16X16X64_I8_vgprcd_e64:
case V_MFMA_F32_16X16X32_F16_e64:
case V_MFMA_F32_16X16X32_F16_vgprcd_e64:
- // Distinction marker to simplify mapping to the programming guide
+ // GFX9 Shader Programming Guide lists those SMFMAC separately, but for
+ // intended purposes here all those instructions are the same. This comment
+ // is to simplify reverse mapping to the SPG.
case V_SMFMAC_F32_16X16X64_BF16_e64:
case V_SMFMAC_I32_16X16X128_I8_e64:
case V_SMFMAC_F32_16X16X128_BF8_BF8_e64:
>From 56474cb9adcc29db1338ebbd410b9c7a9eb00010 Mon Sep 17 00:00:00 2001
From: Alexey Sachkov <alexey.sachkov at amd.com>
Date: Tue, 18 Aug 2026 10:23:26 -0500
Subject: [PATCH 03/11] Add VMEM to coexec rules
---
llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h | 12 ++++++------
1 file changed, 6 insertions(+), 6 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
index 7a7d424d0d3f0..ffcb65c71441d 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
@@ -398,7 +398,7 @@ inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
case V_SMFMAC_F32_16X16X64_F16_e64:
Res.TotalWindow = 8;
AllowCoExec(Res, CoExecMask::SALU, 1);
- AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU, 2);
+ AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU | CoExecMask::VMEM, 2);
AllowCoExec(Res, CoExecMask::WMMA, 4);
return Res;
@@ -415,7 +415,7 @@ inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64:
Res.TotalWindow = 12;
AllowCoExec(Res, CoExecMask::SALU, 1);
- AllowCoExec(Res, CoExecMask::DS, 2);
+ AllowCoExec(Res, CoExecMask::DS | CoExecMask::VMEM, 2);
AllowCoExec(Res, CoExecMask::VALU, 3);
AllowCoExec(Res, CoExecMask::WMMA, 8);
return Res;
@@ -451,7 +451,7 @@ inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
case V_MFMA_F32_32X32X16_F16_vgprcd_e64:
Res.TotalWindow = 8;
AllowCoExec(Res, CoExecMask::SALU, 1);
- AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU, 2);
+ AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU | CoExecMask::VMEM, 2);
AllowCoExec(Res, CoExecMask::WMMA, 4);
return Res;
@@ -478,7 +478,7 @@ inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_vgprcd_e64:
Res.TotalWindow = 20;
AllowCoExec(Res, CoExecMask::SALU, 1);
- AllowCoExec(Res, CoExecMask::DS, 2);
+ AllowCoExec(Res, CoExecMask::DS | CoExecMask::VMEM, 2);
AllowCoExec(Res, CoExecMask::VALU, 3);
AllowCoExec(Res, CoExecMask::WMMA, 16);
return Res;
@@ -493,7 +493,7 @@ inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
case V_SMFMAC_F32_32X32X32_F16_e64:
Res.TotalWindow = 12;
AllowCoExec(Res, CoExecMask::SALU, 1);
- AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU, 4);
+ AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU | CoExecMask::VMEM, 4);
AllowCoExec(Res, CoExecMask::WMMA, 9);
return Res;
@@ -503,7 +503,7 @@ inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
case V_MFMA_F64_16X16X4F64_mac_vgprcd_e64:
case V_MFMA_F64_16X16X4F64_vgprcd_e64:
Res.TotalWindow = 19;
- AllowCoExec(Res, CoExecMask::DS | CoExecMask::SALU, 0);
+ AllowCoExec(Res, CoExecMask::DS | CoExecMask::SALU | CoExecMask::VMEM, 0);
AllowCoExec(Res, CoExecMask::WMMA | CoExecMask::VALU, 18);
return Res;
>From fc3b5c015f8deb5b0103ac6a2dcc1339df9f741a Mon Sep 17 00:00:00 2001
From: Alexey Sachkov <alexey.sachkov at amd.com>
Date: Tue, 18 Aug 2026 10:29:26 -0500
Subject: [PATCH 04/11] Outline AllowCoExec into a separate helper function
instead of a lambda
---
llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h | 19 ++++++++++---------
1 file changed, 10 insertions(+), 9 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
index ffcb65c71441d..d1d9e69ce172f 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
@@ -346,6 +346,16 @@ inline CoExecInfo CoExecInfo::build(unsigned TotalWindow, const char *Pattern) {
return Info;
}
+/// Apply \p ExtraBits to every slot in \p Info starting with \p StartIndex
+/// Used by MFMA co-exec rules, because MFMA co-exec slots are incremental, i.e.
+/// for every slot N it supports all instructions which were supported by the
+/// previous slot N-1 and may support something extra.
+inline void AllowCoExec(CoExecInfo &Info, CoExecMaskT ExtraBits,
+ unsigned StartIndex) {
+ for (unsigned Index = StartIndex; Index < Info.TotalWindow; ++Index)
+ Info.Slots[Index].Mask |= ExtraBits;
+}
+
/// Get co-execution info for a gfx950 MFMA instruction.
/// The occupancy (cycles until the next MFMA may issue) is expressed as the
/// first stage carrying the WMMA bit.
@@ -361,15 +371,6 @@ inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
// needed by printing specific slots in full instead of a key for them.
Res.Pattern = "undefinedundefinedundefinedundefined";
- // MFMA co-exec slots are incremental, i.e. for every slot N it supports all
- // instructions which were supported by the previous slot N-1 and may support
- // something extra.
- auto AllowCoExec = [](CoExecInfo &Info, CoExecMaskT ExtraBits,
- unsigned StartIndex) {
- for (unsigned Index = StartIndex; Index < Info.TotalWindow; ++Index)
- Info.Slots[Index].Mask |= ExtraBits;
- };
-
switch (MI.getOpcode()) {
// 4-cycle occupancy, 8-cycle window.
case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_e64:
>From 66017d0813e0363622fa8c574e0db7f1fb553c99 Mon Sep 17 00:00:00 2001
From: Alexey Sachkov <alexey.sachkov at amd.com>
Date: Thu, 20 Aug 2026 04:23:25 -0500
Subject: [PATCH 05/11] Add unit-test to make sure that all opcodes are covered
---
llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h | 156 ++++++++----------
llvm/unittests/Target/AMDGPU/CMakeLists.txt | 1 +
.../Target/AMDGPU/MFMACoExecRules.cpp | 69 ++++++++
3 files changed, 136 insertions(+), 90 deletions(-)
create mode 100644 llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
index d1d9e69ce172f..d98be317ad3f2 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
@@ -359,7 +359,7 @@ inline void AllowCoExec(CoExecInfo &Info, CoExecMaskT ExtraBits,
/// Get co-execution info for a gfx950 MFMA instruction.
/// The occupancy (cycles until the next MFMA may issue) is expressed as the
/// first stage carrying the WMMA bit.
-inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
+inline CoExecInfo getMFMACoExecInfo(unsigned Opcode) {
CoExecInfo Res;
for (unsigned I = 0; I < MaxCoExecStages; ++I)
Res.Slots[I].Mask = CoExecMask::None;
@@ -371,32 +371,32 @@ inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
// needed by printing specific slots in full instead of a key for them.
Res.Pattern = "undefinedundefinedundefinedundefined";
- switch (MI.getOpcode()) {
+ switch (Opcode) {
// 4-cycle occupancy, 8-cycle window.
- case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f4_f6_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f4_f6_vgprcd_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f6_f4_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f6_f4_vgprcd_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f6_f6_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f6_f6_vgprcd_e64:
- case V_MFMA_F32_16X16X32_BF16_e64:
- case V_MFMA_F32_16X16X32_BF16_vgprcd_e64:
- case V_MFMA_I32_16X16X64_I8_e64:
- case V_MFMA_I32_16X16X64_I8_vgprcd_e64:
- case V_MFMA_F32_16X16X32_F16_e64:
- case V_MFMA_F32_16X16X32_F16_vgprcd_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f6_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f6_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f4_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f4_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f6_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f6_gfx940_vcd:
+ case V_MFMA_F32_16X16X32_BF16_gfx940_acd:
+ case V_MFMA_F32_16X16X32_BF16_gfx940_vcd:
+ case V_MFMA_I32_16X16X64_I8_gfx940_acd:
+ case V_MFMA_I32_16X16X64_I8_gfx940_vcd:
+ case V_MFMA_F32_16X16X32_F16_gfx940_acd:
+ case V_MFMA_F32_16X16X32_F16_gfx940_vcd:
// GFX9 Shader Programming Guide lists those SMFMAC separately, but for
// intended purposes here all those instructions are the same. This comment
// is to simplify reverse mapping to the SPG.
- case V_SMFMAC_F32_16X16X64_BF16_e64:
- case V_SMFMAC_I32_16X16X128_I8_e64:
- case V_SMFMAC_F32_16X16X128_BF8_BF8_e64:
- case V_SMFMAC_F32_16X16X128_BF8_FP8_e64:
- case V_SMFMAC_F32_16X16X128_FP8_BF8_e64:
- case V_SMFMAC_F32_16X16X128_FP8_FP8_e64:
- case V_SMFMAC_F32_16X16X64_F16_e64:
+ case V_SMFMAC_F32_16X16X64_BF16_gfx940:
+ case V_SMFMAC_I32_16X16X128_I8_gfx940:
+ case V_SMFMAC_F32_16X16X128_BF8_BF8_gfx940:
+ case V_SMFMAC_F32_16X16X128_BF8_FP8_gfx940:
+ case V_SMFMAC_F32_16X16X128_FP8_BF8_gfx940:
+ case V_SMFMAC_F32_16X16X128_FP8_FP8_gfx940:
+ case V_SMFMAC_F32_16X16X64_F16_gfx940:
Res.TotalWindow = 8;
AllowCoExec(Res, CoExecMask::SALU, 1);
AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU | CoExecMask::VMEM, 2);
@@ -404,16 +404,16 @@ inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
return Res;
// 8-cycle occupancy, 12-cycle window.
- case V_MFMA_F32_16X16X128_F8F6F4_f4_f8_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f4_f8_vgprcd_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f6_f8_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f6_f8_vgprcd_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f8_f4_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f8_f4_vgprcd_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_vgprcd_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_e64:
- case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f8_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f8_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f8_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f8_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f4_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f4_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_gfx940_vcd:
Res.TotalWindow = 12;
AllowCoExec(Res, CoExecMask::SALU, 1);
AllowCoExec(Res, CoExecMask::DS | CoExecMask::VMEM, 2);
@@ -422,34 +422,20 @@ inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
return Res;
// 4-cycle occupancy, 8-cycle window.
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_mac_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_mac_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_mac_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_mac_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_mac_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_mac_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_mac_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_mac_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_vgprcd_e64:
- case V_MFMA_F32_32X32X16_BF16_e64:
- case V_MFMA_F32_32X32X16_BF16_mac_e64:
- case V_MFMA_F32_32X32X16_BF16_mac_vgprcd_e64:
- case V_MFMA_F32_32X32X16_BF16_vgprcd_e64:
- case V_MFMA_I32_32X32X16I8_e64:
- case V_MFMA_I32_32X32X16I8_mac_e64:
- case V_MFMA_I32_32X32X16I8_mac_vgprcd_e64:
- case V_MFMA_I32_32X32X16I8_vgprcd_e64:
- case V_MFMA_F32_32X32X16_F16_e64:
- case V_MFMA_F32_32X32X16_F16_mac_e64:
- case V_MFMA_F32_32X32X16_F16_mac_vgprcd_e64:
- case V_MFMA_F32_32X32X16_F16_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_gfx940_vcd:
+ case V_MFMA_F32_32X32X16_BF16_gfx940_acd:
+ case V_MFMA_F32_32X32X16_BF16_gfx940_vcd:
+ case V_MFMA_I32_32X32X32_I8_gfx940_acd:
+ case V_MFMA_I32_32X32X32_I8_gfx940_vcd:
+ case V_MFMA_F32_32X32X16_F16_gfx940_acd:
+ case V_MFMA_F32_32X32X16_F16_gfx940_vcd:
Res.TotalWindow = 8;
AllowCoExec(Res, CoExecMask::SALU, 1);
AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU | CoExecMask::VMEM, 2);
@@ -457,26 +443,16 @@ inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
return Res;
// 16-cycle occupancy, 20-cycle window.
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_mac_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_mac_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_mac_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_mac_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_mac_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_mac_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_mac_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_mac_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_mac_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_mac_vgprcd_e64:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_gfx940_vcd:
Res.TotalWindow = 20;
AllowCoExec(Res, CoExecMask::SALU, 1);
AllowCoExec(Res, CoExecMask::DS | CoExecMask::VMEM, 2);
@@ -485,13 +461,13 @@ inline CoExecInfo getMFMACoExecInfo(const MachineInstr &MI) {
return Res;
// 9-cycle occupancy, 12-cycle window.
- case V_SMFMAC_F32_32X32X32_BF16_e64:
- case V_SMFMAC_I32_32X32X64_I8_e64:
- case V_SMFMAC_F32_32X32X64_BF8_BF8_e64:
- case V_SMFMAC_F32_32X32X64_BF8_FP8_e64:
- case V_SMFMAC_F32_32X32X64_FP8_BF8_e64:
- case V_SMFMAC_F32_32X32X64_FP8_FP8_e64:
- case V_SMFMAC_F32_32X32X32_F16_e64:
+ case V_SMFMAC_F32_32X32X32_BF16_gfx940:
+ case V_SMFMAC_I32_32X32X64_I8_gfx940:
+ case V_SMFMAC_F32_32X32X64_BF8_BF8_gfx940:
+ case V_SMFMAC_F32_32X32X64_BF8_FP8_gfx940:
+ case V_SMFMAC_F32_32X32X64_FP8_BF8_gfx940:
+ case V_SMFMAC_F32_32X32X64_FP8_FP8_gfx940:
+ case V_SMFMAC_F32_32X32X32_F16_gfx940:
Res.TotalWindow = 12;
AllowCoExec(Res, CoExecMask::SALU, 1);
AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU | CoExecMask::VMEM, 4);
@@ -520,8 +496,8 @@ inline CoExecInfo getCoExecInfo(const MachineInstr &MI,
const SIInstrInfo &TII) {
unsigned Opc = MI.getOpcode();
- if (TII.isMFMA(MI))
- return getMFMACoExecInfo(MI);
+ if (TII.isMFMA(Opc))
+ return getMFMACoExecInfo(Opc);
// Scaled variants (LD_SCALE rule) absorb the next WMMA in the last I slot.
bool HasScaling = AMDGPU::getHasMatrixScale(Opc);
diff --git a/llvm/unittests/Target/AMDGPU/CMakeLists.txt b/llvm/unittests/Target/AMDGPU/CMakeLists.txt
index 39cced662567d..097f15bd397b9 100644
--- a/llvm/unittests/Target/AMDGPU/CMakeLists.txt
+++ b/llvm/unittests/Target/AMDGPU/CMakeLists.txt
@@ -33,4 +33,5 @@ add_llvm_target_unittest(AMDGPUTests
LiveRegUnits.cpp
PALMetadata.cpp
UniformityAnalysisTest.cpp
+ MFMACoExecRules.cpp
)
diff --git a/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp b/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp
new file mode 100644
index 0000000000000..1321ae4e5c5bb
--- /dev/null
+++ b/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp
@@ -0,0 +1,69 @@
+//===- llvm/unittest/CodeGen/AMDGPUMetadataTest.cpp -----------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+/// \file
+/// Test that MFMA co-exec rules are defined for all gfx950 MFMA instructions.
+//
+//===----------------------------------------------------------------------===//
+
+#include "AMDGPUUnitTests.h"
+#include "gtest/gtest.h"
+
+#include "AMDGPUCoExecInfo.h"
+#include "llvm/MC/MCInstrInfo.h"
+#include "llvm/TargetParser/SubtargetFeature.h"
+
+#include "GCNSubtarget.h"
+#include "SIInstrInfo.h"
+#include "llvm/CodeGen/MachineFunction.h"
+
+#define GET_AVAILABLE_OPCODE_CHECKER
+#include "AMDGPUGenInstrInfo.inc"
+
+using namespace llvm;
+using namespace llvm::AMDGPU;
+
+class AMDGPUMFMACoExecRules : public AMDGPUTestBase {};
+
+TEST_F(AMDGPUMFMACoExecRules, Basic) {
+ std::unique_ptr<GCNTargetMachine> TM =
+ createAMDGPUTargetMachine(Triple("amdgpu-amd-amdhsa"), "gfx950", "");
+ ASSERT_NE(TM, nullptr);
+ auto ST =
+ std::make_unique<GCNSubtarget>(TM->getTargetTriple(), TM->getTargetCPU(),
+ TM->getTargetFeatureString(), *TM);
+ ASSERT_NE(ST, nullptr);
+
+ const MCInstrInfo *MCII = TM->getMCInstrInfo();
+
+ // Get feature bits for gfx950
+ const FeatureBitset &Features950 = ST->getFeatureBits();
+ FeatureBitset Available950 = AMDGPU_MC::computeAvailableFeatures(Features950);
+
+ for (unsigned Op = 0; Op < MCII->getNumOpcodes(); ++Op) {
+ ASSERT_NE(ST->getInstrInfo(), nullptr);
+ if (!ST->getInstrInfo()->isMFMAorWMMA(Op))
+ continue;
+ if (MCII->get(Op).isPseudo())
+ continue;
+
+ FeatureBitset Required = AMDGPU_MC::computeRequiredFeatures(Op);
+ if (!Required.test(AMDGPU_MC::Feature_HasGFX950InstsBit))
+ continue;
+
+ FeatureBitset Missing = (Available950 & Required) ^ Required;
+ bool AvailableOnGfx950 = Missing.none();
+
+ CoExecInfo Info = getMFMACoExecInfo(Op);
+ if (AvailableOnGfx950) {
+ EXPECT_EQ(Info.Slots[0].Mask, CoExecMask::None) << MCII->getName(Op);
+ } else {
+ // Ignore for now/not part of this test
+ }
+ }
+}
>From 7e37d80920700069cb511871bb3710a14b93f0ba Mon Sep 17 00:00:00 2001
From: Alexey Sachkov <alexey.sachkov at amd.com>
Date: Thu, 20 Aug 2026 07:51:41 -0500
Subject: [PATCH 06/11] Add rules for MFMA_LD_SCALE instructions
---
llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h | 36 +++++++++++++++++++++++
1 file changed, 36 insertions(+)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
index d98be317ad3f2..d8099a8cf05e5 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
@@ -387,6 +387,14 @@ inline CoExecInfo getMFMACoExecInfo(unsigned Opcode) {
case V_MFMA_I32_16X16X64_I8_gfx940_vcd:
case V_MFMA_F32_16X16X32_F16_gfx940_acd:
case V_MFMA_F32_16X16X32_F16_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f6_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f6_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f4_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f6_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f6_gfx940_vcd:
// GFX9 Shader Programming Guide lists those SMFMAC separately, but for
// intended purposes here all those instructions are the same. This comment
// is to simplify reverse mapping to the SPG.
@@ -414,6 +422,16 @@ inline CoExecInfo getMFMACoExecInfo(unsigned Opcode) {
case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_gfx940_vcd:
case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_gfx940_acd:
case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f8_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f8_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f4_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f6_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f6_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_gfx940_vcd:
Res.TotalWindow = 12;
AllowCoExec(Res, CoExecMask::SALU, 1);
AllowCoExec(Res, CoExecMask::DS | CoExecMask::VMEM, 2);
@@ -436,6 +454,14 @@ inline CoExecInfo getMFMACoExecInfo(unsigned Opcode) {
case V_MFMA_I32_32X32X32_I8_gfx940_vcd:
case V_MFMA_F32_32X32X16_F16_gfx940_acd:
case V_MFMA_F32_32X32X16_F16_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f4_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f6_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f6_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f4_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f6_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f6_gfx940_vcd:
Res.TotalWindow = 8;
AllowCoExec(Res, CoExecMask::SALU, 1);
AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU | CoExecMask::VMEM, 2);
@@ -453,6 +479,16 @@ inline CoExecInfo getMFMACoExecInfo(unsigned Opcode) {
case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_gfx940_vcd:
case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_gfx940_acd:
case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f8_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f8_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f4_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f6_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f6_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f8_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f8_gfx940_vcd:
Res.TotalWindow = 20;
AllowCoExec(Res, CoExecMask::SALU, 1);
AllowCoExec(Res, CoExecMask::DS | CoExecMask::VMEM, 2);
>From 9e1bb56230d735e97d1b0c3c648d9bf92780d582 Mon Sep 17 00:00:00 2001
From: Alexey Sachkov <alexey.sachkov at amd.com>
Date: Thu, 20 Aug 2026 08:05:24 -0500
Subject: [PATCH 07/11] Apply code review comments
---
llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h | 179 +----------------
.../AMDGPU/AMDGPUCoExecSchedStrategy.cpp | 183 ++++++++++++++++++
2 files changed, 184 insertions(+), 178 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
index d8099a8cf05e5..78b5fb0d60068 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
@@ -346,185 +346,8 @@ inline CoExecInfo CoExecInfo::build(unsigned TotalWindow, const char *Pattern) {
return Info;
}
-/// Apply \p ExtraBits to every slot in \p Info starting with \p StartIndex
-/// Used by MFMA co-exec rules, because MFMA co-exec slots are incremental, i.e.
-/// for every slot N it supports all instructions which were supported by the
-/// previous slot N-1 and may support something extra.
-inline void AllowCoExec(CoExecInfo &Info, CoExecMaskT ExtraBits,
- unsigned StartIndex) {
- for (unsigned Index = StartIndex; Index < Info.TotalWindow; ++Index)
- Info.Slots[Index].Mask |= ExtraBits;
-}
-
/// Get co-execution info for a gfx950 MFMA instruction.
-/// The occupancy (cycles until the next MFMA may issue) is expressed as the
-/// first stage carrying the WMMA bit.
-inline CoExecInfo getMFMACoExecInfo(unsigned Opcode) {
- CoExecInfo Res;
- for (unsigned I = 0; I < MaxCoExecStages; ++I)
- Res.Slots[I].Mask = CoExecMask::None;
-
- // TODO: Implement proper patterns support (for debugging purposes).
- // Existing pattern letters are WMMA-specific and will probably be confusing
- // if used as-is for MFMA. Inventing new MFMA-specific letters is an option,
- // but perhaps the pattern should be instead dynamically reconstructed when
- // needed by printing specific slots in full instead of a key for them.
- Res.Pattern = "undefinedundefinedundefinedundefined";
-
- switch (Opcode) {
- // 4-cycle occupancy, 8-cycle window.
- case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_gfx940_acd:
- case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_gfx940_vcd:
- case V_MFMA_F32_16X16X128_F8F6F4_f4_f6_gfx940_acd:
- case V_MFMA_F32_16X16X128_F8F6F4_f4_f6_gfx940_vcd:
- case V_MFMA_F32_16X16X128_F8F6F4_f6_f4_gfx940_acd:
- case V_MFMA_F32_16X16X128_F8F6F4_f6_f4_gfx940_vcd:
- case V_MFMA_F32_16X16X128_F8F6F4_f6_f6_gfx940_acd:
- case V_MFMA_F32_16X16X128_F8F6F4_f6_f6_gfx940_vcd:
- case V_MFMA_F32_16X16X32_BF16_gfx940_acd:
- case V_MFMA_F32_16X16X32_BF16_gfx940_vcd:
- case V_MFMA_I32_16X16X64_I8_gfx940_acd:
- case V_MFMA_I32_16X16X64_I8_gfx940_vcd:
- case V_MFMA_F32_16X16X32_F16_gfx940_acd:
- case V_MFMA_F32_16X16X32_F16_gfx940_vcd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_gfx940_acd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_gfx940_vcd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f6_gfx940_acd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f6_gfx940_vcd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f4_gfx940_acd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f4_gfx940_vcd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f6_gfx940_acd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f6_gfx940_vcd:
- // GFX9 Shader Programming Guide lists those SMFMAC separately, but for
- // intended purposes here all those instructions are the same. This comment
- // is to simplify reverse mapping to the SPG.
- case V_SMFMAC_F32_16X16X64_BF16_gfx940:
- case V_SMFMAC_I32_16X16X128_I8_gfx940:
- case V_SMFMAC_F32_16X16X128_BF8_BF8_gfx940:
- case V_SMFMAC_F32_16X16X128_BF8_FP8_gfx940:
- case V_SMFMAC_F32_16X16X128_FP8_BF8_gfx940:
- case V_SMFMAC_F32_16X16X128_FP8_FP8_gfx940:
- case V_SMFMAC_F32_16X16X64_F16_gfx940:
- Res.TotalWindow = 8;
- AllowCoExec(Res, CoExecMask::SALU, 1);
- AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU | CoExecMask::VMEM, 2);
- AllowCoExec(Res, CoExecMask::WMMA, 4);
- return Res;
-
- // 8-cycle occupancy, 12-cycle window.
- case V_MFMA_F32_16X16X128_F8F6F4_f4_f8_gfx940_acd:
- case V_MFMA_F32_16X16X128_F8F6F4_f4_f8_gfx940_vcd:
- case V_MFMA_F32_16X16X128_F8F6F4_f6_f8_gfx940_acd:
- case V_MFMA_F32_16X16X128_F8F6F4_f6_f8_gfx940_vcd:
- case V_MFMA_F32_16X16X128_F8F6F4_f8_f4_gfx940_acd:
- case V_MFMA_F32_16X16X128_F8F6F4_f8_f4_gfx940_vcd:
- case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_gfx940_acd:
- case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_gfx940_vcd:
- case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_gfx940_acd:
- case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_gfx940_vcd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f8_gfx940_acd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f8_gfx940_vcd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f8_gfx940_acd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f8_gfx940_vcd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f4_gfx940_acd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f4_gfx940_vcd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f6_gfx940_acd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f6_gfx940_vcd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_gfx940_acd:
- case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_gfx940_vcd:
- Res.TotalWindow = 12;
- AllowCoExec(Res, CoExecMask::SALU, 1);
- AllowCoExec(Res, CoExecMask::DS | CoExecMask::VMEM, 2);
- AllowCoExec(Res, CoExecMask::VALU, 3);
- AllowCoExec(Res, CoExecMask::WMMA, 8);
- return Res;
-
- // 4-cycle occupancy, 8-cycle window.
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_gfx940_acd:
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_gfx940_vcd:
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_gfx940_acd:
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_gfx940_vcd:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_gfx940_acd:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_gfx940_vcd:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_gfx940_acd:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_gfx940_vcd:
- case V_MFMA_F32_32X32X16_BF16_gfx940_acd:
- case V_MFMA_F32_32X32X16_BF16_gfx940_vcd:
- case V_MFMA_I32_32X32X32_I8_gfx940_acd:
- case V_MFMA_I32_32X32X32_I8_gfx940_vcd:
- case V_MFMA_F32_32X32X16_F16_gfx940_acd:
- case V_MFMA_F32_32X32X16_F16_gfx940_vcd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f4_gfx940_acd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f4_gfx940_vcd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f6_gfx940_acd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f6_gfx940_vcd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f4_gfx940_acd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f4_gfx940_vcd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f6_gfx940_acd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f6_gfx940_vcd:
- Res.TotalWindow = 8;
- AllowCoExec(Res, CoExecMask::SALU, 1);
- AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU | CoExecMask::VMEM, 2);
- AllowCoExec(Res, CoExecMask::WMMA, 4);
- return Res;
-
- // 16-cycle occupancy, 20-cycle window.
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_gfx940_acd:
- case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_gfx940_vcd:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_gfx940_acd:
- case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_gfx940_vcd:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_gfx940_acd:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_gfx940_vcd:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_gfx940_acd:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_gfx940_vcd:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_gfx940_acd:
- case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_gfx940_vcd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f8_gfx940_acd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f8_gfx940_vcd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f8_gfx940_acd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f8_gfx940_vcd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f4_gfx940_acd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f4_gfx940_vcd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f6_gfx940_acd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f6_gfx940_vcd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f8_gfx940_acd:
- case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f8_gfx940_vcd:
- Res.TotalWindow = 20;
- AllowCoExec(Res, CoExecMask::SALU, 1);
- AllowCoExec(Res, CoExecMask::DS | CoExecMask::VMEM, 2);
- AllowCoExec(Res, CoExecMask::VALU, 3);
- AllowCoExec(Res, CoExecMask::WMMA, 16);
- return Res;
-
- // 9-cycle occupancy, 12-cycle window.
- case V_SMFMAC_F32_32X32X32_BF16_gfx940:
- case V_SMFMAC_I32_32X32X64_I8_gfx940:
- case V_SMFMAC_F32_32X32X64_BF8_BF8_gfx940:
- case V_SMFMAC_F32_32X32X64_BF8_FP8_gfx940:
- case V_SMFMAC_F32_32X32X64_FP8_BF8_gfx940:
- case V_SMFMAC_F32_32X32X64_FP8_FP8_gfx940:
- case V_SMFMAC_F32_32X32X32_F16_gfx940:
- Res.TotalWindow = 12;
- AllowCoExec(Res, CoExecMask::SALU, 1);
- AllowCoExec(Res, CoExecMask::DS | CoExecMask::VALU | CoExecMask::VMEM, 4);
- AllowCoExec(Res, CoExecMask::WMMA, 9);
- return Res;
-
- // 18-cycle occupancy, 19-cycle window.
- case V_MFMA_F64_16X16X4F64_e64:
- case V_MFMA_F64_16X16X4F64_mac_e64:
- case V_MFMA_F64_16X16X4F64_mac_vgprcd_e64:
- case V_MFMA_F64_16X16X4F64_vgprcd_e64:
- Res.TotalWindow = 19;
- AllowCoExec(Res, CoExecMask::DS | CoExecMask::SALU | CoExecMask::VMEM, 0);
- AllowCoExec(Res, CoExecMask::WMMA | CoExecMask::VALU, 18);
- return Res;
-
- default:
- // Default fallback: permissive 8-cycle pattern
- return CoExecInfo::build(9, "AAAAAAAAA");
- }
-}
+CoExecInfo getMFMACoExecInfo(unsigned Opcode);
/// Get co-execution info for a WMMA instruction, selecting the per-cycle slot
/// pattern from the opcode (and operand formats for the F8F6F4 variants).
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
index c68e1b8508148..61162d5df8931 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
@@ -44,6 +44,189 @@ static SUnit *pickOnlyChoice(SchedBoundary &Zone) {
return OnlyChoice;
}
+/// Apply \p ExtraBits to every slot in \p Info starting with \p StartIndex
+/// Used by MFMA co-exec rules, because MFMA co-exec slots are incremental, i.e.
+/// for every slot N it supports all instructions which were supported by the
+/// previous slot N-1 and may support something extra.
+static void allowCoExec(llvm::AMDGPU::CoExecInfo &Info,
+ llvm::AMDGPU::CoExecMaskT ExtraBits,
+ unsigned StartIndex) {
+ for (unsigned Index = StartIndex; Index < Info.TotalWindow; ++Index)
+ Info.Slots[Index].Mask |= ExtraBits;
+}
+
+/// Get co-execution info for a gfx950 MFMA instruction.
+/// The occupancy (cycles until the next MFMA may issue) is expressed as the
+/// first stage carrying the WMMA bit.
+llvm::AMDGPU::CoExecInfo llvm::AMDGPU::getMFMACoExecInfo(unsigned Opcode) {
+ using namespace llvm;
+ using namespace llvm::AMDGPU;
+ CoExecInfo Res;
+ for (unsigned I = 0; I < MaxCoExecStages; ++I)
+ Res.Slots[I].Mask = CoExecMask::None;
+
+ // TODO: Implement proper patterns support (for debugging purposes).
+ // Existing pattern letters are WMMA-specific and will probably be confusing
+ // if used as-is for MFMA. Inventing new MFMA-specific letters is an option,
+ // but perhaps the pattern should be instead dynamically reconstructed when
+ // needed by printing specific slots in full instead of a key for them.
+ Res.Pattern = "undefinedundefinedundefinedundefined";
+
+ switch (Opcode) {
+ // 4-cycle occupancy, 8-cycle window.
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f6_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f6_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f4_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f4_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f6_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f6_gfx940_vcd:
+ case V_MFMA_F32_16X16X32_BF16_gfx940_acd:
+ case V_MFMA_F32_16X16X32_BF16_gfx940_vcd:
+ case V_MFMA_I32_16X16X64_I8_gfx940_acd:
+ case V_MFMA_I32_16X16X64_I8_gfx940_vcd:
+ case V_MFMA_F32_16X16X32_F16_gfx940_acd:
+ case V_MFMA_F32_16X16X32_F16_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f6_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f6_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f4_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f6_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f6_gfx940_vcd:
+ // GFX9 Shader Programming Guide lists those SMFMAC separately, but for
+ // intended purposes here all those instructions are the same. This comment
+ // is to simplify reverse mapping to the SPG.
+ case V_SMFMAC_F32_16X16X64_BF16_gfx940:
+ case V_SMFMAC_I32_16X16X128_I8_gfx940:
+ case V_SMFMAC_F32_16X16X128_BF8_BF8_gfx940:
+ case V_SMFMAC_F32_16X16X128_BF8_FP8_gfx940:
+ case V_SMFMAC_F32_16X16X128_FP8_BF8_gfx940:
+ case V_SMFMAC_F32_16X16X128_FP8_FP8_gfx940:
+ case V_SMFMAC_F32_16X16X64_F16_gfx940:
+ Res.TotalWindow = 8;
+ allowCoExec(Res, CoExecMask::SALU, 1);
+ allowCoExec(Res, CoExecMask::DS | CoExecMask::VALU | CoExecMask::VMEM, 2);
+ allowCoExec(Res, CoExecMask::WMMA, 4);
+ return Res;
+
+ // 8-cycle occupancy, 12-cycle window.
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f8_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f8_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f8_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f8_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f4_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f4_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_gfx940_acd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f8_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f8_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f4_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f6_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f6_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_gfx940_acd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_gfx940_vcd:
+ Res.TotalWindow = 12;
+ allowCoExec(Res, CoExecMask::SALU, 1);
+ allowCoExec(Res, CoExecMask::DS | CoExecMask::VMEM, 2);
+ allowCoExec(Res, CoExecMask::VALU, 3);
+ allowCoExec(Res, CoExecMask::WMMA, 8);
+ return Res;
+
+ // 4-cycle occupancy, 8-cycle window.
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_gfx940_vcd:
+ case V_MFMA_F32_32X32X16_BF16_gfx940_acd:
+ case V_MFMA_F32_32X32X16_BF16_gfx940_vcd:
+ case V_MFMA_I32_32X32X32_I8_gfx940_acd:
+ case V_MFMA_I32_32X32X32_I8_gfx940_vcd:
+ case V_MFMA_F32_32X32X16_F16_gfx940_acd:
+ case V_MFMA_F32_32X32X16_F16_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f4_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f6_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f6_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f4_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f6_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f6_gfx940_vcd:
+ Res.TotalWindow = 8;
+ allowCoExec(Res, CoExecMask::SALU, 1);
+ allowCoExec(Res, CoExecMask::DS | CoExecMask::VALU | CoExecMask::VMEM, 2);
+ allowCoExec(Res, CoExecMask::WMMA, 4);
+ return Res;
+
+ // 16-cycle occupancy, 20-cycle window.
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_gfx940_acd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f8_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f8_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f4_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f6_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f6_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f8_gfx940_acd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f8_gfx940_vcd:
+ Res.TotalWindow = 20;
+ allowCoExec(Res, CoExecMask::SALU, 1);
+ allowCoExec(Res, CoExecMask::DS | CoExecMask::VMEM, 2);
+ allowCoExec(Res, CoExecMask::VALU, 3);
+ allowCoExec(Res, CoExecMask::WMMA, 16);
+ return Res;
+
+ // 9-cycle occupancy, 12-cycle window.
+ case V_SMFMAC_F32_32X32X32_BF16_gfx940:
+ case V_SMFMAC_I32_32X32X64_I8_gfx940:
+ case V_SMFMAC_F32_32X32X64_BF8_BF8_gfx940:
+ case V_SMFMAC_F32_32X32X64_BF8_FP8_gfx940:
+ case V_SMFMAC_F32_32X32X64_FP8_BF8_gfx940:
+ case V_SMFMAC_F32_32X32X64_FP8_FP8_gfx940:
+ case V_SMFMAC_F32_32X32X32_F16_gfx940:
+ Res.TotalWindow = 12;
+ allowCoExec(Res, CoExecMask::SALU, 1);
+ allowCoExec(Res, CoExecMask::DS | CoExecMask::VALU | CoExecMask::VMEM, 4);
+ allowCoExec(Res, CoExecMask::WMMA, 9);
+ return Res;
+
+ // 18-cycle occupancy, 19-cycle window.
+ case V_MFMA_F64_16X16X4F64_e64:
+ case V_MFMA_F64_16X16X4F64_mac_e64:
+ case V_MFMA_F64_16X16X4F64_mac_vgprcd_e64:
+ case V_MFMA_F64_16X16X4F64_vgprcd_e64:
+ Res.TotalWindow = 19;
+ allowCoExec(Res, CoExecMask::DS | CoExecMask::SALU | CoExecMask::VMEM, 0);
+ allowCoExec(Res, CoExecMask::WMMA | CoExecMask::VALU, 18);
+ return Res;
+
+ default:
+ // Default fallback: permissive 8-cycle pattern
+ return CoExecInfo::build(9, "AAAAAAAAA");
+ }
+}
+
InstructionFlavor llvm::AMDGPU::classifyFlavor(const MachineInstr &MI,
const SIInstrInfo &SII) {
if (MI.isDebugInstr())
>From 939e5dfcc0d9669b3bb84cef012354894480ccae Mon Sep 17 00:00:00 2001
From: Alexey Sachkov <alexey.sachkov at amd.com>
Date: Thu, 20 Aug 2026 08:42:44 -0500
Subject: [PATCH 08/11] Make sure to cover all MFMA opcodes, including
pseudo-instructions
---
.../AMDGPU/AMDGPUCoExecSchedStrategy.cpp | 140 ++++++++++++++++++
.../Target/AMDGPU/MFMACoExecRules.cpp | 2 -
2 files changed, 140 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
index 61162d5df8931..9d5ef3a4b49e2 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
@@ -74,37 +74,66 @@ llvm::AMDGPU::CoExecInfo llvm::AMDGPU::getMFMACoExecInfo(unsigned Opcode) {
switch (Opcode) {
// 4-cycle occupancy, 8-cycle window.
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64:
case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_gfx940_acd:
case V_MFMA_F32_16X16X128_F8F6F4_f4_f4_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f6_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f6_vgprcd_e64:
case V_MFMA_F32_16X16X128_F8F6F4_f4_f6_gfx940_acd:
case V_MFMA_F32_16X16X128_F8F6F4_f4_f6_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f4_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f4_vgprcd_e64:
case V_MFMA_F32_16X16X128_F8F6F4_f6_f4_gfx940_acd:
case V_MFMA_F32_16X16X128_F8F6F4_f6_f4_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f6_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f6_vgprcd_e64:
case V_MFMA_F32_16X16X128_F8F6F4_f6_f6_gfx940_acd:
case V_MFMA_F32_16X16X128_F8F6F4_f6_f6_gfx940_vcd:
+ case V_MFMA_F32_16X16X32_BF16_e64:
+ case V_MFMA_F32_16X16X32_BF16_vgprcd_e64:
case V_MFMA_F32_16X16X32_BF16_gfx940_acd:
case V_MFMA_F32_16X16X32_BF16_gfx940_vcd:
+ case V_MFMA_I32_16X16X64_I8_e64:
+ case V_MFMA_I32_16X16X64_I8_vgprcd_e64:
case V_MFMA_I32_16X16X64_I8_gfx940_acd:
case V_MFMA_I32_16X16X64_I8_gfx940_vcd:
+ case V_MFMA_F32_16X16X32_F16_e64:
+ case V_MFMA_F32_16X16X32_F16_vgprcd_e64:
case V_MFMA_F32_16X16X32_F16_gfx940_acd:
case V_MFMA_F32_16X16X32_F16_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_gfx940_acd:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f6_e64:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f6_vgprcd_e64:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f6_gfx940_acd:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f6_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f4_e64:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f4_vgprcd_e64:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f4_gfx940_acd:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f6_e64:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f6_vgprcd_e64:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f6_gfx940_acd:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f6_gfx940_vcd:
// GFX9 Shader Programming Guide lists those SMFMAC separately, but for
// intended purposes here all those instructions are the same. This comment
// is to simplify reverse mapping to the SPG.
+ case V_SMFMAC_F32_16X16X64_BF16_e64:
case V_SMFMAC_F32_16X16X64_BF16_gfx940:
+ case V_SMFMAC_I32_16X16X128_I8_e64:
case V_SMFMAC_I32_16X16X128_I8_gfx940:
+ case V_SMFMAC_F32_16X16X128_BF8_BF8_e64:
case V_SMFMAC_F32_16X16X128_BF8_BF8_gfx940:
+ case V_SMFMAC_F32_16X16X128_BF8_FP8_e64:
case V_SMFMAC_F32_16X16X128_BF8_FP8_gfx940:
+ case V_SMFMAC_F32_16X16X128_FP8_BF8_e64:
case V_SMFMAC_F32_16X16X128_FP8_BF8_gfx940:
+ case V_SMFMAC_F32_16X16X128_FP8_FP8_e64:
case V_SMFMAC_F32_16X16X128_FP8_FP8_gfx940:
+ case V_SMFMAC_F32_16X16X64_F16_e64:
case V_SMFMAC_F32_16X16X64_F16_gfx940:
Res.TotalWindow = 8;
allowCoExec(Res, CoExecMask::SALU, 1);
@@ -113,24 +142,44 @@ llvm::AMDGPU::CoExecInfo llvm::AMDGPU::getMFMACoExecInfo(unsigned Opcode) {
return Res;
// 8-cycle occupancy, 12-cycle window.
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f8_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f4_f8_vgprcd_e64:
case V_MFMA_F32_16X16X128_F8F6F4_f4_f8_gfx940_acd:
case V_MFMA_F32_16X16X128_F8F6F4_f4_f8_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f8_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f6_f8_vgprcd_e64:
case V_MFMA_F32_16X16X128_F8F6F4_f6_f8_gfx940_acd:
case V_MFMA_F32_16X16X128_F8F6F4_f6_f8_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f4_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f4_vgprcd_e64:
case V_MFMA_F32_16X16X128_F8F6F4_f8_f4_gfx940_acd:
case V_MFMA_F32_16X16X128_F8F6F4_f8_f4_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_vgprcd_e64:
case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_gfx940_acd:
case V_MFMA_F32_16X16X128_F8F6F4_f8_f6_gfx940_vcd:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_e64:
+ case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64:
case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_gfx940_acd:
case V_MFMA_F32_16X16X128_F8F6F4_f8_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f8_e64:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f8_vgprcd_e64:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f8_gfx940_acd:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f8_e64:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f8_vgprcd_e64:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f8_gfx940_acd:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f6_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f4_e64:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f4_vgprcd_e64:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f4_gfx940_acd:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f4_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f6_e64:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f6_vgprcd_e64:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f6_gfx940_acd:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f6_gfx940_vcd:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_e64:
+ case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_gfx940_acd:
case V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_gfx940_vcd:
Res.TotalWindow = 12;
@@ -141,18 +190,46 @@ llvm::AMDGPU::CoExecInfo llvm::AMDGPU::getMFMACoExecInfo(unsigned Opcode) {
return Res;
// 4-cycle occupancy, 8-cycle window.
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_vgprcd_e64:
case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_gfx940_acd:
case V_MFMA_F32_32X32X64_F8F6F4_f4_f4_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_vgprcd_e64:
case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_gfx940_acd:
case V_MFMA_F32_32X32X64_F8F6F4_f4_f6_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_vgprcd_e64:
case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_gfx940_acd:
case V_MFMA_F32_32X32X64_F8F6F4_f6_f4_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_vgprcd_e64:
case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_gfx940_acd:
case V_MFMA_F32_32X32X64_F8F6F4_f6_f6_gfx940_vcd:
+ case V_MFMA_F32_32X32X16_BF16_e64:
+ case V_MFMA_F32_32X32X16_BF16_mac_e64:
+ case V_MFMA_F32_32X32X16_BF16_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X16_BF16_vgprcd_e64:
case V_MFMA_F32_32X32X16_BF16_gfx940_acd:
case V_MFMA_F32_32X32X16_BF16_gfx940_vcd:
+ case V_MFMA_I32_32X32X32_I8_e64:
+ case V_MFMA_I32_32X32X32_I8_mac_e64:
+ case V_MFMA_I32_32X32X32_I8_mac_vgprcd_e64:
+ case V_MFMA_I32_32X32X32_I8_vgprcd_e64:
case V_MFMA_I32_32X32X32_I8_gfx940_acd:
case V_MFMA_I32_32X32X32_I8_gfx940_vcd:
+ case V_MFMA_F32_32X32X16_F16_e64:
+ case V_MFMA_F32_32X32X16_F16_mac_e64:
+ case V_MFMA_F32_32X32X16_F16_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X16_F16_vgprcd_e64:
case V_MFMA_F32_32X32X16_F16_gfx940_acd:
case V_MFMA_F32_32X32X16_F16_gfx940_vcd:
case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f4_gfx940_acd:
@@ -163,6 +240,22 @@ llvm::AMDGPU::CoExecInfo llvm::AMDGPU::getMFMACoExecInfo(unsigned Opcode) {
case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f4_gfx940_vcd:
case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f6_gfx940_acd:
case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f6_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f4_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f6_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f4_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f6_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f4_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f6_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f4_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f6_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f4_mac_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f6_mac_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f4_mac_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f6_mac_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f4_mac_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f6_mac_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f4_mac_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f6_mac_vgprcd_e64:
Res.TotalWindow = 8;
allowCoExec(Res, CoExecMask::SALU, 1);
allowCoExec(Res, CoExecMask::DS | CoExecMask::VALU | CoExecMask::VMEM, 2);
@@ -170,16 +263,56 @@ llvm::AMDGPU::CoExecInfo llvm::AMDGPU::getMFMACoExecInfo(unsigned Opcode) {
return Res;
// 16-cycle occupancy, 20-cycle window.
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_vgprcd_e64:
case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_gfx940_acd:
case V_MFMA_F32_32X32X64_F8F6F4_f4_f8_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_vgprcd_e64:
case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_gfx940_acd:
case V_MFMA_F32_32X32X64_F8F6F4_f6_f8_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_vgprcd_e64:
case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_gfx940_acd:
case V_MFMA_F32_32X32X64_F8F6F4_f8_f4_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_vgprcd_e64:
case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_gfx940_acd:
case V_MFMA_F32_32X32X64_F8F6F4_f8_f6_gfx940_vcd:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_mac_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_mac_vgprcd_e64:
+ case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_vgprcd_e64:
case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_gfx940_acd:
case V_MFMA_F32_32X32X64_F8F6F4_f8_f8_gfx940_vcd:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f8_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f8_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f4_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f6_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f8_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f8_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f8_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f4_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f6_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f8_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f8_mac_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f8_mac_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f4_mac_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f6_mac_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f8_mac_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f8_mac_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f8_mac_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f4_mac_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f6_mac_vgprcd_e64:
+ case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f8_f8_mac_vgprcd_e64:
case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f8_gfx940_acd:
case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f4_f8_gfx940_vcd:
case V_MFMA_SCALE_F32_32X32X64_F8F6F4_f6_f8_gfx940_acd:
@@ -198,12 +331,19 @@ llvm::AMDGPU::CoExecInfo llvm::AMDGPU::getMFMACoExecInfo(unsigned Opcode) {
return Res;
// 9-cycle occupancy, 12-cycle window.
+ case V_SMFMAC_F32_32X32X32_BF16_e64:
case V_SMFMAC_F32_32X32X32_BF16_gfx940:
+ case V_SMFMAC_I32_32X32X64_I8_e64:
case V_SMFMAC_I32_32X32X64_I8_gfx940:
+ case V_SMFMAC_F32_32X32X64_BF8_BF8_e64:
case V_SMFMAC_F32_32X32X64_BF8_BF8_gfx940:
+ case V_SMFMAC_F32_32X32X64_BF8_FP8_e64:
case V_SMFMAC_F32_32X32X64_BF8_FP8_gfx940:
+ case V_SMFMAC_F32_32X32X64_FP8_BF8_e64:
case V_SMFMAC_F32_32X32X64_FP8_BF8_gfx940:
+ case V_SMFMAC_F32_32X32X64_FP8_FP8_e64:
case V_SMFMAC_F32_32X32X64_FP8_FP8_gfx940:
+ case V_SMFMAC_F32_32X32X32_F16_e64:
case V_SMFMAC_F32_32X32X32_F16_gfx940:
Res.TotalWindow = 12;
allowCoExec(Res, CoExecMask::SALU, 1);
diff --git a/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp b/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp
index 1321ae4e5c5bb..5b48d0e029c3d 100644
--- a/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp
+++ b/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp
@@ -49,8 +49,6 @@ TEST_F(AMDGPUMFMACoExecRules, Basic) {
ASSERT_NE(ST->getInstrInfo(), nullptr);
if (!ST->getInstrInfo()->isMFMAorWMMA(Op))
continue;
- if (MCII->get(Op).isPseudo())
- continue;
FeatureBitset Required = AMDGPU_MC::computeRequiredFeatures(Op);
if (!Required.test(AMDGPU_MC::Feature_HasGFX950InstsBit))
>From 40bdb6a2ea53b2c790af9b4af404781470b91387 Mon Sep 17 00:00:00 2001
From: Alexey Sachkov <alexey.sachkov at amd.com>
Date: Thu, 20 Aug 2026 09:56:02 -0500
Subject: [PATCH 09/11] Cleanup and clarify the new unit-test
---
.../Target/AMDGPU/MFMACoExecRules.cpp | 23 ++++++-------------
1 file changed, 7 insertions(+), 16 deletions(-)
diff --git a/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp b/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp
index 5b48d0e029c3d..c588b768d8332 100644
--- a/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp
+++ b/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp
@@ -30,38 +30,29 @@ using namespace llvm::AMDGPU;
class AMDGPUMFMACoExecRules : public AMDGPUTestBase {};
-TEST_F(AMDGPUMFMACoExecRules, Basic) {
+TEST_F(AMDGPUMFMACoExecRules, GFX950) {
std::unique_ptr<GCNTargetMachine> TM =
createAMDGPUTargetMachine(Triple("amdgpu-amd-amdhsa"), "gfx950", "");
- ASSERT_NE(TM, nullptr);
auto ST =
std::make_unique<GCNSubtarget>(TM->getTargetTriple(), TM->getTargetCPU(),
TM->getTargetFeatureString(), *TM);
- ASSERT_NE(ST, nullptr);
const MCInstrInfo *MCII = TM->getMCInstrInfo();
- // Get feature bits for gfx950
- const FeatureBitset &Features950 = ST->getFeatureBits();
- FeatureBitset Available950 = AMDGPU_MC::computeAvailableFeatures(Features950);
-
for (unsigned Op = 0; Op < MCII->getNumOpcodes(); ++Op) {
- ASSERT_NE(ST->getInstrInfo(), nullptr);
if (!ST->getInstrInfo()->isMFMAorWMMA(Op))
continue;
+ // Filter out pre-gfx950 MFMA instructions
FeatureBitset Required = AMDGPU_MC::computeRequiredFeatures(Op);
if (!Required.test(AMDGPU_MC::Feature_HasGFX950InstsBit))
continue;
- FeatureBitset Missing = (Available950 & Required) ^ Required;
- bool AvailableOnGfx950 = Missing.none();
-
CoExecInfo Info = getMFMACoExecInfo(Op);
- if (AvailableOnGfx950) {
- EXPECT_EQ(Info.Slots[0].Mask, CoExecMask::None) << MCII->getName(Op);
- } else {
- // Ignore for now/not part of this test
- }
+ // None of gfx950 MFMAs can co-execute anything in their issue slot and our
+ // fallback for unrecognized MFMA instructions is to allow everything in
+ // every slot, so here we check that we did not encounter the fallback
+ // definition of co-exec rules.
+ EXPECT_EQ(Info.Slots[0].Mask, CoExecMask::None) << MCII->getName(Op);
}
}
>From f80cc730cc11412f0c0dc98ff38cd0c4258cd1ed Mon Sep 17 00:00:00 2001
From: Vlad <Vladislav.Romanov at amd.com>
Date: Tue, 25 Aug 2026 02:31:50 -0500
Subject: [PATCH 10/11] add llvm ir test
---
.../AMDGPU/coexec-mfma-interleave-gfx950.ll | 137 ++++++++++++++++++
1 file changed, 137 insertions(+)
create mode 100644 llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll
diff --git a/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll b/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll
new file mode 100644
index 0000000000000..8cd93a4388e1f
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll
@@ -0,0 +1,137 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -amdgpu-sched-strategy=coexec %s -o - 2>&1 | FileCheck %s
+
+; Test MFMA coexec window interleaving on gfx950 from LLVM IR.
+; Three independent MFMAs with independent VALUs and SALUs should be
+; interleaved: 2 VALUs + 1 SALU per MFMA shadow.
+
+declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half>, <8 x half>, <4 x float>, i32 immarg, i32 immarg, i32 immarg)
+
+define amdgpu_kernel void @mfma_16x16_interleave(
+; CHECK-LABEL: mfma_16x16_interleave:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x0
+; CHECK-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x8
+; CHECK-NEXT: v_and_b32_e32 v2, 0x3ff, v0
+; CHECK-NEXT: v_mov_b32_e32 v1, 0
+; CHECK-NEXT: v_lshlrev_b32_e32 v0, 4, v2
+; CHECK-NEXT: v_mul_hi_i32_i24_e32 v15, -12, v2
+; CHECK-NEXT: v_mul_i32_i24_e32 v14, -12, v2
+; CHECK-NEXT: s_waitcnt lgkmcnt(0)
+; CHECK-NEXT: v_lshl_add_u64 v[16:17], s[0:1], 0, v[0:1]
+; CHECK-NEXT: global_load_dwordx4 v[2:5], v0, s[0:1]
+; CHECK-NEXT: global_load_dwordx4 v[6:9], v0, s[0:1] offset:16
+; CHECK-NEXT: global_load_dwordx4 v[10:13], v0, s[0:1] offset:32
+; CHECK-NEXT: v_lshl_add_u64 v[18:19], v[16:17], 0, v[14:15]
+; CHECK-NEXT: global_load_dwordx4 v[14:17], v[18:19], off
+; CHECK-NEXT: global_load_dwordx2 v[30:31], v[18:19], off offset:16
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: ; sched_barrier mask(0x00000000)
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[18:21], v[2:5], v[6:9], v[2:5]
+; CHECK-NEXT: s_add_i32 s0, s8, s9
+; CHECK-NEXT: v_add_u32_e32 v22, v14, v15
+; CHECK-NEXT: v_add_u32_e32 v23, v15, v16
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[26:29], v[2:5], v[6:9], v[6:9]
+; CHECK-NEXT: s_add_i32 s1, s10, s11
+; CHECK-NEXT: v_add_u32_e32 v24, v16, v17
+; CHECK-NEXT: v_add_u32_e32 v25, v17, v30
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[2:5], v[2:5], v[6:9], v[10:13]
+; CHECK-NEXT: s_add_i32 s2, s12, s13
+; CHECK-NEXT: v_add_u32_e32 v6, v30, v31
+; CHECK-NEXT: v_add_u32_e32 v7, v31, v14
+; CHECK-NEXT: ; sched_barrier mask(0x00000000)
+; CHECK-NEXT: v_mov_b32_e32 v8, s0
+; CHECK-NEXT: v_mov_b32_e32 v9, s1
+; CHECK-NEXT: v_mov_b32_e32 v0, s2
+; CHECK-NEXT: global_store_dwordx4 v1, v[18:21], s[14:15]
+; CHECK-NEXT: global_store_dwordx4 v1, v[26:29], s[14:15] offset:16
+; CHECK-NEXT: global_store_dwordx4 v1, v[2:5], s[14:15] offset:32
+; CHECK-NEXT: global_store_dwordx4 v1, v[22:25], s[14:15] offset:48
+; CHECK-NEXT: global_store_dwordx4 v1, v[6:9], s[14:15] offset:64
+; CHECK-NEXT: global_store_dword v1, v0, s[14:15] offset:80
+; CHECK-NEXT: s_endpgm
+ ptr addrspace(1) %ptr,
+ i32 %s0, i32 %s1, i32 %s2, i32 %s3,
+ i32 %s4, i32 %s5,
+ ptr addrspace(1) %out) {
+
+ ; Load all operands using divergent (per-lane) addresses to ensure VGPRs.
+ %tid = call i32 @llvm.amdgcn.workitem.id.x()
+
+ ; Load MFMA src and acc operands.
+ %mbase = getelementptr <8 x half>, ptr addrspace(1) %ptr, i32 %tid
+ %src0 = load <8 x half>, ptr addrspace(1) %mbase
+ %ms1 = getelementptr <8 x half>, ptr addrspace(1) %mbase, i32 1
+ %src1 = load <8 x half>, ptr addrspace(1) %ms1
+ %abase = getelementptr <4 x float>, ptr addrspace(1) %ptr, i32 %tid
+ %acc0 = load <4 x float>, ptr addrspace(1) %abase
+ %ap1 = getelementptr <4 x float>, ptr addrspace(1) %abase, i32 1
+ %acc1 = load <4 x float>, ptr addrspace(1) %ap1
+ %ap2 = getelementptr <4 x float>, ptr addrspace(1) %abase, i32 2
+ %acc2 = load <4 x float>, ptr addrspace(1) %ap2
+
+ ; Load VALU operands.
+ %vbase = getelementptr i32, ptr addrspace(1) %ptr, i32 %tid
+ %p1 = getelementptr i32, ptr addrspace(1) %vbase, i32 1
+ %p2 = getelementptr i32, ptr addrspace(1) %vbase, i32 2
+ %p3 = getelementptr i32, ptr addrspace(1) %vbase, i32 3
+ %p4 = getelementptr i32, ptr addrspace(1) %vbase, i32 4
+ %p5 = getelementptr i32, ptr addrspace(1) %vbase, i32 5
+ %v0 = load i32, ptr addrspace(1) %vbase
+ %v1 = load i32, ptr addrspace(1) %p1
+ %v2 = load i32, ptr addrspace(1) %p2
+ %v3 = load i32, ptr addrspace(1) %p3
+ %v4 = load i32, ptr addrspace(1) %p4
+ %v5 = load i32, ptr addrspace(1) %p5
+
+ ; Wait for all loads to complete before MFMAs.
+ call void @llvm.amdgcn.s.waitcnt(i32 0)
+ call void @llvm.amdgcn.sched.barrier(i32 0)
+
+ ; Three independent MFMAs (reuse src0/src1, different acc).
+ %m0 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %src0, <8 x half> %src1, <4 x float> %acc0, i32 0, i32 0, i32 0)
+ %m1 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %src0, <8 x half> %src1, <4 x float> %acc1, i32 0, i32 0, i32 0)
+ %m2 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %src0, <8 x half> %src1, <4 x float> %acc2, i32 0, i32 0, i32 0)
+
+ ; Independent VALUs (divergent i32 adds → v_add_u32)
+ %a0 = add i32 %v0, %v1
+ %a1 = add i32 %v1, %v2
+ %a2 = add i32 %v2, %v3
+ %a3 = add i32 %v3, %v4
+ %a4 = add i32 %v4, %v5
+ %a5 = add i32 %v5, %v0
+
+ ; Independent SALUs
+ %sa0 = add i32 %s0, %s1
+ %sa1 = add i32 %s2, %s3
+ %sa2 = add i32 %s4, %s5
+
+ ; Scheduling barrier to separate compute from stores.
+ call void @llvm.amdgcn.sched.barrier(i32 0)
+
+ ; Store everything to prevent DCE.
+ %out1 = getelementptr <4 x float>, ptr addrspace(1) %out, i32 1
+ %out2 = getelementptr <4 x float>, ptr addrspace(1) %out, i32 2
+ store <4 x float> %m0, ptr addrspace(1) %out
+ store <4 x float> %m1, ptr addrspace(1) %out1
+ store <4 x float> %m2, ptr addrspace(1) %out2
+ %outa = getelementptr i32, ptr addrspace(1) %out, i32 12
+ store i32 %a0, ptr addrspace(1) %outa
+ %outa1 = getelementptr i32, ptr addrspace(1) %outa, i32 1
+ store i32 %a1, ptr addrspace(1) %outa1
+ %outa2 = getelementptr i32, ptr addrspace(1) %outa, i32 2
+ store i32 %a2, ptr addrspace(1) %outa2
+ %outa3 = getelementptr i32, ptr addrspace(1) %outa, i32 3
+ store i32 %a3, ptr addrspace(1) %outa3
+ %outa4 = getelementptr i32, ptr addrspace(1) %outa, i32 4
+ store i32 %a4, ptr addrspace(1) %outa4
+ %outa5 = getelementptr i32, ptr addrspace(1) %outa, i32 5
+ store i32 %a5, ptr addrspace(1) %outa5
+ %outs = getelementptr i32, ptr addrspace(1) %out, i32 18
+ store i32 %sa0, ptr addrspace(1) %outs
+ %outs1 = getelementptr i32, ptr addrspace(1) %outs, i32 1
+ store i32 %sa1, ptr addrspace(1) %outs1
+ %outs2 = getelementptr i32, ptr addrspace(1) %outs, i32 2
+ store i32 %sa2, ptr addrspace(1) %outs2
+ ret void
+}
>From 39d0775a4b868d20367a1f28b3470790d31f3349 Mon Sep 17 00:00:00 2001
From: Vlad <Vladislav.Romanov at amd.com>
Date: Tue, 25 Aug 2026 04:39:23 -0500
Subject: [PATCH 11/11] updat test
---
llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll b/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll
index 8cd93a4388e1f..c290874a5d89b 100644
--- a/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll
+++ b/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll
@@ -1,5 +1,5 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -amdgpu-sched-strategy=coexec %s -o - 2>&1 | FileCheck %s
+; RUN: llc -mtriple=amdgpu9.50-amd-amdhsa -amdgpu-sched-strategy=coexec %s -o - 2>&1 | FileCheck %s
; Test MFMA coexec window interleaving on gfx950 from LLVM IR.
; Three independent MFMAs with independent VALUs and SALUs should be
@@ -23,8 +23,8 @@ define amdgpu_kernel void @mfma_16x16_interleave(
; CHECK-NEXT: global_load_dwordx4 v[6:9], v0, s[0:1] offset:16
; CHECK-NEXT: global_load_dwordx4 v[10:13], v0, s[0:1] offset:32
; CHECK-NEXT: v_lshl_add_u64 v[18:19], v[16:17], 0, v[14:15]
-; CHECK-NEXT: global_load_dwordx4 v[14:17], v[18:19], off
; CHECK-NEXT: global_load_dwordx2 v[30:31], v[18:19], off offset:16
+; CHECK-NEXT: global_load_dwordx4 v[14:17], v[18:19], off
; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; CHECK-NEXT: ; sched_barrier mask(0x00000000)
; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[18:21], v[2:5], v[6:9], v[2:5]
More information about the llvm-commits
mailing list