[llvm] [AMDGPU] Add initial version of gfx950 MFMA co-exec rules (PR #214682)
Alexey Sachkov via llvm-commits
llvm-commits at lists.llvm.org
Thu Aug 20 02:24:12 PDT 2026
https://github.com/AlexeySachkov 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 1/4] [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 2/4] 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 3/4] 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 4/4] 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:
More information about the llvm-commits
mailing list