[llvm-branch-commits] [llvm] [AMDGPU] Fix the packed F32 unpack window in SIPreEmitPeephole (PR #233721)
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Oct 19 23:59:54 PDT 2026
https://github.com/lijinpei-amd created https://github.com/llvm/llvm-project/pull/233721
collectUnpackingCandidates unpacked a `v_pk_*_f32` only if its two halves
finished before NumPasses - 1, i.e. before the MFMA completes. A packed F32
cannot issue until NumPasses + 2 after an XDL (NumPasses + 3 after an XDL2x),
though, so a packed instruction placed between those points still stalled.
Unpacking pays off whenever the pair finishes before that cycle, so use
getPackedVALUIssueCycle - getVALUIssueCycle + 1 as the window.
Also:
- Skip SGEMM and DGEMM. They hold every VALU until NumPasses + 2, so
unpacking after them only adds an issue slot.
- Charge S_NOP its wait states when counting cycles. The larger window
would otherwise unpack instructions that issue long after the MFMA.
getVALUIssueCycle and getPackedVALUIssueCycle are new. They return the first
cycle, in 4-clock issue slots counted from the MFMA's issue, at which a VALU,
or a packed VALU or DOT, can issue after an MFMA on gfx942 and gfx950. The
rules fit ATT thread traces taken on MI308X (gfx942, with the peak TOPS
limiter enabled) and MI350X (gfx950) for every MFMA family:
- An XDL holds VALU off until 1 + (widest A/B source in bytes) / 16.
- PK_F32 (`v_pk_{add,fma,mul}_f32`, `v_pk_mov_b32`) and DOT issue at
NumPasses + 2, one cycle more for XDL2x (an A source of 16 bytes or more).
Packed 16-bit ops issue like any other VALU.
- SGEMM and DGEMM hold every VALU until NumPasses + 2.
A scaled MFMA follows the same rules when it has to wait for an earlier MFMA
to free the pipe, as back-to-back MFMAs in a loop do: its scale load overlaps
that wait. Issued into an idle pipe, it starts one cycle late, which delays
VALU, PK_F32, DOT and the next MFMA by one cycle. The helpers model the
back-to-back case. Streams of scaled MFMAs run at the unscaled rate.
DOT includes the VOP2 `v_dot*c` forms, which are not IsPacked, so
getPackedVALUIssueCycle accepts any packed or DOT opcode. A unit test checks
every measured MFMA and probe combination.
Timed with s_memtime on MI308X (peak TOPS limiter enabled) and MI350X, in
clocks per MFMA followed by four `v_pk_fma_f32` (six for F8F6F4):
- XDL 16x16x16: 36 -> 32
- BF16 XDL2x 16x16x32: 40 -> 36
- F8F6F4 XDL2x 16x16x128: 56 -> 52
- SGEMM 16x16x4: 68 -> 56
- DGEMM 16x16x4: 68 -> 56 (MI308X), 104 -> 88 (MI350X)
Packed 16-bit instructions and DOTs are never unpacked, so they don't change:
after an XDL 16x16x16, four `v_pk_add_f16` take 20 clocks per MFMA and four
`v_dot2_f32_f16` take 40.
With the limiter disabled, MI308X never co-issues VALU with an XDL. The XDL
case becomes 36 -> 40, and four `v_pk_add_f16` take 32.
Assisted-by: Claude Opus 5.5
>From 0a75fd37d8398e7626f4dd3fb4082a339860a8ce Mon Sep 17 00:00:00 2001
From: Li Jinpei <jinpli at amd.com>
Date: Mon, 19 Oct 2026 20:42:56 +0800
Subject: [PATCH 1/3] [AMDGPU] Add MFMA VALU and packed-VALU issue cycle
helpers
Add getVALUIssueCycle and getPackedVALUIssueCycle. They return the first
cycle, in 4-clock issue slots counted from the MFMA's issue, at which a VALU,
or a packed VALU or DOT, can issue after an MFMA on gfx942 and gfx950.
The rules fit ATT thread traces taken on MI308X (gfx942, with the peak TOPS
limiter enabled) and MI350X (gfx950) for every MFMA family:
- An XDL holds VALU off until 1 + (widest A/B source in bytes) / 16.
- PK_F32 (v_pk_{add,fma,mul}_f32, v_pk_mov_b32) and DOT issue at
NumPasses + 2, one cycle more for XDL2x (an A source of 16 bytes or more).
Packed 16-bit ops issue like any other VALU.
- SGEMM and DGEMM hold every VALU until NumPasses + 2.
A scaled MFMA follows the same rules when it has to wait for an earlier
MFMA to free the pipe, as back-to-back MFMAs in a loop do: its scale load
overlaps that wait. Issued into an idle pipe, it starts one cycle late, which
delays VALU, PK_F32, DOT and the next MFMA by one cycle. The helpers model
the back-to-back case. Streams of scaled MFMAs run at the unscaled rate.
DOT includes the VOP2 v_dot*c forms, which are not IsPacked, so
getPackedVALUIssueCycle accepts any packed or DOT opcode.
The unit test checks every measured MFMA and probe combination.
Assisted-by: Claude Opus 5.5
---
llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h | 12 ++
.../AMDGPU/AMDGPUCoExecSchedStrategy.cpp | 52 ++++++
.../Target/AMDGPU/MFMACoExecRules.cpp | 164 +++++++++++++++++-
3 files changed, 226 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
index ab1a7525e91ce..ae26260a50838 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h
@@ -371,6 +371,18 @@ inline CoExecInfo CoExecInfo::build(unsigned UnitOccupancy,
/// Get co-execution info for a gfx950 MFMA instruction.
CoExecInfo getMFMACoExecInfo(unsigned Opcode);
+/// Get the first cycle, counted from \p MFMA issuing at cycle 0, at which a
+/// VALU can issue on gfx942 with the peak TOPS limiter enabled, or on gfx950.
+/// A scaled MFMA is assumed to issue behind another MFMA; issued into an idle
+/// MFMA pipe, it takes one more cycle.
+unsigned getVALUIssueCycle(const MachineInstr &MFMA, const SIInstrInfo &TII);
+
+/// Same as getVALUIssueCycle, for a packed VALU or a DOT with opcode
+/// \p PackedVALUOpc.
+unsigned getPackedVALUIssueCycle(const MachineInstr &MFMA,
+ unsigned PackedVALUOpc,
+ const SIInstrInfo &TII);
+
/// 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,
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
index 0a4961285584d..2555afadf1262 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
@@ -385,6 +385,58 @@ llvm::AMDGPU::CoExecInfo llvm::AMDGPU::getMFMACoExecInfo(unsigned Opcode) {
}
}
+static unsigned getMFMANumPasses(const MachineInstr &MFMA,
+ const SIInstrInfo &TII) {
+ unsigned NumPasses = TII.getRepeatRate(MFMA);
+ assert(NumPasses >= 2 && "MFMA has no pass count in the scheduling model");
+ return NumPasses;
+}
+
+static unsigned getMFMASrcSize(const MachineInstr &MFMA, AMDGPU::OpName Name,
+ const SIInstrInfo &TII) {
+ unsigned Opc = MFMA.getOpcode();
+ return TII.getOpSize(Opc, AMDGPU::getNamedOperandIdx(Opc, Name));
+}
+
+unsigned llvm::AMDGPU::getVALUIssueCycle(const MachineInstr &MFMA,
+ const SIInstrInfo &TII) {
+ assert(TII.getSubtarget().hasGFX940Insts() && SIInstrInfo::isMFMA(MFMA) &&
+ "expected a gfx942 or gfx950 MFMA");
+ if (!TII.isXDL(MFMA))
+ return getMFMANumPasses(MFMA, TII) + 2;
+
+ // An XDL holds VALU off for one cycle per 16 bytes of its widest A/B source.
+ unsigned MaxSrcSize = std::max(getMFMASrcSize(MFMA, OpName::src0, TII),
+ getMFMASrcSize(MFMA, OpName::src1, TII));
+ return 1 + MaxSrcSize / 16;
+}
+
+unsigned llvm::AMDGPU::getPackedVALUIssueCycle(const MachineInstr &MFMA,
+ unsigned PackedVALUOpc,
+ const SIInstrInfo &TII) {
+ assert(TII.getSubtarget().hasGFX940Insts() && SIInstrInfo::isMFMA(MFMA) &&
+ (TII.isPacked(PackedVALUOpc) || TII.isDOT(PackedVALUOpc)) &&
+ "expected a gfx942 or gfx950 MFMA and a packed VALU or DOT");
+ switch (PackedVALUOpc) {
+ case AMDGPU::V_PK_ADD_F32:
+ case AMDGPU::V_PK_FMA_F32:
+ case AMDGPU::V_PK_MOV_B32:
+ case AMDGPU::V_PK_MUL_F32:
+ break;
+ default:
+ if (!TII.isDOT(PackedVALUOpc))
+ return getVALUIssueCycle(MFMA, TII);
+ }
+
+ unsigned NumPasses = getMFMANumPasses(MFMA, TII);
+ if (!TII.isXDL(MFMA))
+ return NumPasses + 2;
+
+ // XDL2x (gfx950 MFMAs with a 16-byte or wider A) adds another cycle.
+ bool IsXDL2x = getMFMASrcSize(MFMA, OpName::src0, TII) >= 16;
+ return NumPasses + 2 + IsXDL2x;
+}
+
InstructionFlavor llvm::AMDGPU::classifyFlavor(const MachineInstr &MI,
const SIInstrInfo &SII) {
if (MI.isDebugInstr())
diff --git a/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp b/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp
index c588b768d8332..55395552f720a 100644
--- a/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp
+++ b/llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp
@@ -1,4 +1,4 @@
-//===- llvm/unittest/CodeGen/AMDGPUMetadataTest.cpp -----------------------===//
+//===- llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp -------------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
@@ -7,7 +7,7 @@
//===----------------------------------------------------------------------===//
//
/// \file
-/// Test that MFMA co-exec rules are defined for all gfx950 MFMA instructions.
+/// Test MFMA co-exec rules.
//
//===----------------------------------------------------------------------===//
@@ -21,6 +21,7 @@
#include "GCNSubtarget.h"
#include "SIInstrInfo.h"
#include "llvm/CodeGen/MachineFunction.h"
+#include "llvm/CodeGen/MachineModuleInfo.h"
#define GET_AVAILABLE_OPCODE_CHECKER
#include "AMDGPUGenInstrInfo.inc"
@@ -56,3 +57,162 @@ TEST_F(AMDGPUMFMACoExecRules, GFX950) {
EXPECT_EQ(Info.Slots[0].Mask, CoExecMask::None) << MCII->getName(Op);
}
}
+
+namespace {
+struct IssueCycles {
+ unsigned MFMAOpc;
+ int64_t CBSZ, BLGP;
+ unsigned VALU, PackedF32;
+};
+
+class MFMAIssueCycleTest {
+public:
+ explicit MFMAIssueCycleTest(std::unique_ptr<GCNTargetMachine> TheTM)
+ : TM(std::move(TheTM)),
+ ST(TM->getTargetTriple(), std::string(TM->getTargetCPU()),
+ std::string(TM->getTargetFeatureString()), *TM),
+ Mod("Module", Ctx), MMI(TM.get()) {
+ Mod.setDataLayout(TM->getTargetTriple().computeDataLayout());
+ auto *F = Function::Create(FunctionType::get(Type::getVoidTy(Ctx), false),
+ GlobalValue::ExternalLinkage, "Test", &Mod);
+ MF = std::make_unique<MachineFunction>(*F, *TM, ST, MMI.getContext(), 0);
+ }
+
+ // Builds an MFMA with placeholder operands; CBSZ/BLGP select F8F6F4 formats.
+ const MachineInstr &buildMFMA(unsigned Opc, int64_t CBSZ, int64_t BLGP) {
+ const SIInstrInfo &TII = *ST.getInstrInfo();
+ const MCInstrDesc &Desc = TII.get(Opc);
+ MachineInstr *MI = MF->CreateMachineInstr(Desc, DebugLoc());
+ for (auto [Idx, Info] : enumerate(Desc.operands())) {
+ if (Info.RegClass == -1)
+ MI->addOperand(*MF, MachineOperand::CreateImm(0));
+ else
+ MI->addOperand(*MF, MachineOperand::CreateReg(Register(), Idx < 1));
+ }
+ if (MachineOperand *Op = TII.getNamedOperand(*MI, AMDGPU::OpName::cbsz))
+ Op->setImm(CBSZ);
+ if (MachineOperand *Op = TII.getNamedOperand(*MI, AMDGPU::OpName::blgp))
+ Op->setImm(BLGP);
+ return *MI;
+ }
+
+ void check(ArrayRef<IssueCycles> Cases) {
+ const SIInstrInfo &TII = *ST.getInstrInfo();
+ for (const IssueCycles &C : Cases) {
+ const MachineInstr &MFMA = buildMFMA(C.MFMAOpc, C.CBSZ, C.BLGP);
+ StringRef Name = TII.getName(C.MFMAOpc);
+ EXPECT_EQ(getVALUIssueCycle(MFMA, TII), C.VALU) << Name;
+ EXPECT_EQ(getPackedVALUIssueCycle(MFMA, V_PK_ADD_F16, TII), C.VALU)
+ << Name;
+ EXPECT_EQ(getPackedVALUIssueCycle(MFMA, V_PK_FMA_F32, TII), C.PackedF32)
+ << Name;
+ EXPECT_EQ(getPackedVALUIssueCycle(MFMA, V_DOT2_F32_F16, TII), C.PackedF32)
+ << Name;
+ }
+ }
+
+ void checkPacked(unsigned MFMAOpc, ArrayRef<unsigned> PackedOpcs,
+ unsigned Expected) {
+ const SIInstrInfo &TII = *ST.getInstrInfo();
+ const MachineInstr &MFMA = buildMFMA(MFMAOpc, 0, 0);
+ for (unsigned PK : PackedOpcs)
+ EXPECT_EQ(getPackedVALUIssueCycle(MFMA, PK, TII), Expected)
+ << TII.getName(PK);
+ }
+
+private:
+ std::unique_ptr<GCNTargetMachine> TM;
+ GCNSubtarget ST;
+ LLVMContext Ctx;
+ Module Mod;
+ MachineModuleInfo MMI;
+ std::unique_ptr<MachineFunction> MF;
+};
+} // namespace
+
+// Expected cycles were measured with ATT on MI308X with the peak TOPS limiter
+// enabled and on MI350X; one cycle is 4 clocks. Scaled MFMAs were measured
+// issuing right behind another MFMA.
+TEST_F(AMDGPUMFMACoExecRules, GFX942IssueCycles) {
+ auto TM =
+ createAMDGPUTargetMachine(Triple("amdgcn-amd-amdhsa"), "gfx942", "");
+ if (!TM)
+ GTEST_SKIP();
+ MFMAIssueCycleTest Test(std::move(TM));
+ Test.checkPacked(V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64,
+ {V_PK_MUL_F32, V_PK_ADD_F32, V_PK_MOV_B32, V_DOT2_I32_I16},
+ 6);
+ Test.checkPacked(V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64,
+ {V_PK_FMA_F16, V_PK_MUL_F16, V_PK_ADD_U16, V_PK_MUL_LO_U16},
+ 1);
+ Test.checkPacked(V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64,
+ {V_DOT2C_F32_F16_e32, V_DOT4C_I32_I8_e32}, 6);
+ Test.checkPacked(V_MFMA_F32_32X32X8BF16_1K_vgprcd_e64,
+ {V_DOT2C_F32_F16_e32, V_DOT4C_I32_I8_e32}, 10);
+ Test.checkPacked(V_MFMA_F32_16X16X4F32_vgprcd_e64,
+ {V_DOT2C_F32_F16_e32, V_DOT4C_I32_I8_e32}, 10);
+ Test.check({
+ {V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64, 0, 0, 1, 6},
+ {V_MFMA_F32_16X16X16BF16_1K_e64, 0, 0, 1, 6},
+ {V_MFMA_F32_4X4X4BF16_1K_vgprcd_e64, 0, 0, 1, 4},
+ {V_MFMA_F32_32X32X8BF16_1K_vgprcd_e64, 0, 0, 1, 10},
+ {V_MFMA_F32_16X16X16F16_vgprcd_e64, 0, 0, 1, 6},
+ {V_MFMA_I32_16X16X32I8_vgprcd_e64, 0, 0, 1, 6},
+ {V_MFMA_F32_16X16X32_FP8_FP8_vgprcd_e64, 0, 0, 1, 6},
+ {V_MFMA_F32_16X16X8XF32_vgprcd_e64, 0, 0, 1, 6},
+ {V_SMFMAC_F32_16X16X32_F16_e64, 0, 0, 2, 6},
+ {V_MFMA_F32_16X16X4F32_vgprcd_e64, 0, 0, 10, 10},
+ {V_MFMA_F64_16X16X4F64_vgprcd_e64, 0, 0, 10, 10},
+ {V_MFMA_F64_4X4X4F64_vgprcd_e64, 0, 0, 6, 6},
+ });
+}
+
+TEST_F(AMDGPUMFMACoExecRules, GFX950IssueCycles) {
+ auto TM =
+ createAMDGPUTargetMachine(Triple("amdgcn-amd-amdhsa"), "gfx950", "");
+ if (!TM)
+ GTEST_SKIP();
+ MFMAIssueCycleTest Test(std::move(TM));
+ Test.checkPacked(
+ V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64,
+ {V_DOT2C_F32_F16_e32, V_DOT4C_I32_I8_e32, V_DOT2C_F32_BF16_e32}, 6);
+ Test.checkPacked(V_MFMA_F32_32X32X8BF16_1K_vgprcd_e64,
+ {V_DOT2C_F32_F16_e32, V_DOT4C_I32_I8_e32}, 10);
+ Test.checkPacked(V_MFMA_F32_16X16X4F32_vgprcd_e64,
+ {V_DOT2C_F32_F16_e32, V_DOT4C_I32_I8_e32}, 10);
+ Test.checkPacked(
+ V_MFMA_F32_16X16X32_BF16_vgprcd_e64,
+ {V_DOT2C_F32_F16_e32, V_DOT4C_I32_I8_e32, V_DOT2C_F32_BF16_e32}, 7);
+ Test.checkPacked(V_MFMA_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64,
+ {V_DOT2C_F32_F16_e32, V_DOT4C_I32_I8_e32}, 11);
+ Test.checkPacked(V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64,
+ {V_DOT2C_F32_F16_e32, V_DOT4C_I32_I8_e32}, 11);
+ Test.check({
+ {V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64, 0, 0, 1, 6},
+ {V_MFMA_F32_4X4X4BF16_1K_vgprcd_e64, 0, 0, 1, 4},
+ {V_MFMA_F32_32X32X8BF16_1K_vgprcd_e64, 0, 0, 1, 10},
+ {V_SMFMAC_F32_16X16X32_F16_e64, 0, 0, 2, 6},
+ {V_MFMA_F32_16X16X32_BF16_vgprcd_e64, 0, 0, 2, 7},
+ {V_MFMA_F32_16X16X32_BF16_e64, 0, 0, 2, 7},
+ {V_MFMA_F32_16X16X32_F16_vgprcd_e64, 0, 0, 2, 7},
+ {V_MFMA_I32_16X16X64_I8_vgprcd_e64, 0, 0, 2, 7},
+ {V_MFMA_F32_32X32X16_BF16_vgprcd_e64, 0, 0, 2, 11},
+ {V_MFMA_F32_32X32X16_F16_vgprcd_e64, 0, 0, 2, 11},
+ {V_MFMA_I32_32X32X32_I8_vgprcd_e64, 0, 0, 2, 11},
+ {V_MFMA_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64, 4, 4, 2, 7},
+ {V_MFMA_F32_16X16X128_F8F6F4_f6_f6_vgprcd_e64, 2, 2, 2, 7},
+ {V_MFMA_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64, 0, 0, 3, 11},
+ {V_MFMA_F32_16X16X128_F8F6F4_f8_f4_vgprcd_e64, 0, 4, 3, 11},
+ {V_MFMA_F32_32X32X64_F8F6F4_f4_f4_vgprcd_e64, 4, 4, 2, 11},
+ {V_MFMA_F32_32X32X64_F8F6F4_f8_f8_vgprcd_e64, 0, 0, 3, 19},
+ {V_MFMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64, 0, 0, 3, 11},
+ {V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64, 4, 4, 2, 7},
+ {V_SMFMAC_F32_16X16X64_BF16_e64, 0, 0, 3, 7},
+ {V_SMFMAC_F32_32X32X32_BF16_e64, 0, 0, 3, 11},
+ {V_SMFMAC_F32_16X16X128_FP8_FP8_e64, 0, 0, 3, 7},
+ {V_SMFMAC_F32_32X32X64_FP8_FP8_e64, 0, 0, 3, 11},
+ {V_MFMA_F32_16X16X4F32_vgprcd_e64, 0, 0, 10, 10},
+ {V_MFMA_F64_16X16X4F64_vgprcd_e64, 0, 0, 18, 18},
+ {V_MFMA_F64_4X4X4F64_vgprcd_e64, 0, 0, 6, 6},
+ });
+}
>From 466ca54156cbb2bb62e33f11a2fd77c3e1be09cf Mon Sep 17 00:00:00 2001
From: Li Jinpei <jinpli at amd.com>
Date: Mon, 19 Oct 2026 22:29:05 +0800
Subject: [PATCH 2/3] [AMDGPU] Precommit SIPreEmitPeephole unpack window tests
(NFC)
Add cases where packed F32 instructions follow a legacy XDL, a BF16 XDL2x,
an F8F6F4 XDL2x (whose 32-byte A source holds VALU off one cycle longer),
an SGEMM, a DGEMM, and an XDL followed by long S_NOPs. Also add cases where
packed 16-bit instructions or DOTs, which the pass never unpacks, follow a
legacy XDL. The CHECK lines record what the pass does today.
Assisted-by: Claude Opus 5.5
---
...coissue-insts-post-ra-scheduler-gfx950.mir | 55 ++++
...ck-non-coissue-insts-post-ra-scheduler.mir | 271 ++++++++++++++++++
2 files changed, 326 insertions(+)
create mode 100644 llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler-gfx950.mir
diff --git a/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler-gfx950.mir b/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler-gfx950.mir
new file mode 100644
index 0000000000000..96313d997526b
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler-gfx950.mir
@@ -0,0 +1,55 @@
+# 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=si-pre-emit-peephole -o - %s | FileCheck -check-prefix=GFX950 %s
+
+---
+name: test_unpack_window_xdl2x
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr20_vgpr21_vgpr22_vgpr23, $vgpr24_vgpr25_vgpr26_vgpr27, $vgpr8_vgpr9
+ ; GFX950-LABEL: name: test_unpack_window_xdl2x
+ ; GFX950: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr20_vgpr21_vgpr22_vgpr23, $vgpr24_vgpr25_vgpr26_vgpr27, $vgpr8_vgpr9
+ ; GFX950-NEXT: {{ $}}
+ ; GFX950-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X32_BF16_vgprcd_e64 $vgpr20_vgpr21_vgpr22_vgpr23, $vgpr24_vgpr25_vgpr26_vgpr27, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: S_ENDPGM 0
+ $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X32_BF16_vgprcd_e64 $vgpr20_vgpr21_vgpr22_vgpr23, $vgpr24_vgpr25_vgpr26_vgpr27, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ S_ENDPGM 0
+...
+---
+name: test_unpack_window_xdl2x_f8
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31, $vgpr32_vgpr33_vgpr34_vgpr35_vgpr36_vgpr37_vgpr38_vgpr39, $vgpr8_vgpr9
+ ; GFX950-LABEL: name: test_unpack_window_xdl2x_f8
+ ; GFX950: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31, $vgpr32_vgpr33_vgpr34_vgpr35_vgpr36_vgpr37_vgpr38_vgpr39, $vgpr8_vgpr9
+ ; GFX950-NEXT: {{ $}}
+ ; GFX950-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64 $vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31, $vgpr32_vgpr33_vgpr34_vgpr35_vgpr36_vgpr37_vgpr38_vgpr39, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr12 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr13 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr14 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr15 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr18_vgpr19 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr20_vgpr21 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: S_ENDPGM 0
+ $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X128_F8F6F4_f8_f8_vgprcd_e64 $vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31, $vgpr32_vgpr33_vgpr34_vgpr35_vgpr36_vgpr37_vgpr38_vgpr39, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, implicit $mode, implicit $exec
+ $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr18_vgpr19 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr20_vgpr21 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ S_ENDPGM 0
+...
diff --git a/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler.mir b/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler.mir
index 7cd6f57e0036c..b56933d627329 100644
--- a/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler.mir
+++ b/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler.mir
@@ -1167,3 +1167,274 @@ body: |
$vgpr5 = V_MOV_B32_e32 killed $sgpr15, implicit $exec, implicit $exec
renamable $vgpr16_vgpr17 = nofpexcept V_PK_MUL_F32 8, killed $sgpr30_sgpr31, 11, killed $vgpr4_vgpr5, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
S_ENDPGM 0
+...
+---
+name: test_unpack_window_xdl
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8_vgpr9
+ ; GFX950-LABEL: name: test_unpack_window_xdl
+ ; GFX950: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8_vgpr9
+ ; GFX950-NEXT: {{ $}}
+ ; GFX950-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: S_ENDPGM 0
+ ;
+ ; GFX942-LABEL: name: test_unpack_window_xdl
+ ; GFX942: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8_vgpr9
+ ; GFX942-NEXT: {{ $}}
+ ; GFX942-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: S_ENDPGM 0
+ ;
+ ; GFX90A-LABEL: name: test_unpack_window_xdl
+ ; GFX90A: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8_vgpr9
+ ; GFX90A-NEXT: {{ $}}
+ ; GFX90A-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: S_ENDPGM 0
+ $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ S_ENDPGM 0
+...
+---
+name: test_unpack_window_counts_s_nop
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8_vgpr9
+ ; GFX950-LABEL: name: test_unpack_window_counts_s_nop
+ ; GFX950: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8_vgpr9
+ ; GFX950-NEXT: {{ $}}
+ ; GFX950-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: S_NOP 7
+ ; GFX950-NEXT: S_NOP 7
+ ; GFX950-NEXT: $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: S_ENDPGM 0
+ ;
+ ; GFX942-LABEL: name: test_unpack_window_counts_s_nop
+ ; GFX942: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8_vgpr9
+ ; GFX942-NEXT: {{ $}}
+ ; GFX942-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: S_NOP 7
+ ; GFX942-NEXT: S_NOP 7
+ ; GFX942-NEXT: $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: S_ENDPGM 0
+ ;
+ ; GFX90A-LABEL: name: test_unpack_window_counts_s_nop
+ ; GFX90A: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8_vgpr9
+ ; GFX90A-NEXT: {{ $}}
+ ; GFX90A-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: S_NOP 7
+ ; GFX90A-NEXT: S_NOP 7
+ ; GFX90A-NEXT: $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: S_ENDPGM 0
+ $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ S_NOP 7
+ S_NOP 7
+ $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ S_ENDPGM 0
+...
+---
+name: test_no_unpack_after_dgemm
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7, $vgpr30_vgpr31, $vgpr8_vgpr9
+ ; GFX950-LABEL: name: test_no_unpack_after_dgemm
+ ; GFX950: liveins: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7, $vgpr30_vgpr31, $vgpr8_vgpr9
+ ; GFX950-NEXT: {{ $}}
+ ; GFX950-NEXT: early-clobber $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7 = V_MFMA_F64_16X16X4F64_vgprcd_e64 $vgpr30_vgpr31, $vgpr30_vgpr31, $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr12 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr13 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr14 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr15 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr16 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr17 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: S_ENDPGM 0
+ ;
+ ; GFX942-LABEL: name: test_no_unpack_after_dgemm
+ ; GFX942: liveins: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7, $vgpr30_vgpr31, $vgpr8_vgpr9
+ ; GFX942-NEXT: {{ $}}
+ ; GFX942-NEXT: early-clobber $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7 = V_MFMA_F64_16X16X4F64_vgprcd_e64 $vgpr30_vgpr31, $vgpr30_vgpr31, $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr12 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr13 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr14 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr15 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: S_ENDPGM 0
+ ;
+ ; GFX90A-LABEL: name: test_no_unpack_after_dgemm
+ ; GFX90A: liveins: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7, $vgpr30_vgpr31, $vgpr8_vgpr9
+ ; GFX90A-NEXT: {{ $}}
+ ; GFX90A-NEXT: early-clobber $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7 = V_MFMA_F64_16X16X4F64_vgprcd_e64 $vgpr30_vgpr31, $vgpr30_vgpr31, $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: S_ENDPGM 0
+ $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7 = V_MFMA_F64_16X16X4F64_vgprcd_e64 $vgpr30_vgpr31, $vgpr30_vgpr31, $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ S_ENDPGM 0
+...
+
+---
+name: test_no_unpack_after_sgemm
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr30, $vgpr31, $vgpr8_vgpr9
+ ; GFX950-LABEL: name: test_no_unpack_after_sgemm
+ ; GFX950: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr30, $vgpr31, $vgpr8_vgpr9
+ ; GFX950-NEXT: {{ $}}
+ ; GFX950-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X4F32_vgprcd_e64 $vgpr30, $vgpr31, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr12 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr13 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr14 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr15 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: S_ENDPGM 0
+ ;
+ ; GFX942-LABEL: name: test_no_unpack_after_sgemm
+ ; GFX942: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr30, $vgpr31, $vgpr8_vgpr9
+ ; GFX942-NEXT: {{ $}}
+ ; GFX942-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X4F32_vgprcd_e64 $vgpr30, $vgpr31, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr12 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr13 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr14 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr15 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: S_ENDPGM 0
+ ;
+ ; GFX90A-LABEL: name: test_no_unpack_after_sgemm
+ ; GFX90A: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr30, $vgpr31, $vgpr8_vgpr9
+ ; GFX90A-NEXT: {{ $}}
+ ; GFX90A-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X4F32_vgprcd_e64 $vgpr30, $vgpr31, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: S_ENDPGM 0
+ $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X4F32_vgprcd_e64 $vgpr30, $vgpr31, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ S_ENDPGM 0
+...
+
+---
+name: test_no_unpack_pk16_after_xdl
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8, $vgpr9
+ ; GFX950-LABEL: name: test_no_unpack_pk16_after_xdl
+ ; GFX950: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8, $vgpr9
+ ; GFX950-NEXT: {{ $}}
+ ; GFX950-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr10 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr11 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr12 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr13 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: S_ENDPGM 0
+ ;
+ ; GFX942-LABEL: name: test_no_unpack_pk16_after_xdl
+ ; GFX942: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8, $vgpr9
+ ; GFX942-NEXT: {{ $}}
+ ; GFX942-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr10 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr11 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr12 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr13 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: S_ENDPGM 0
+ ;
+ ; GFX90A-LABEL: name: test_no_unpack_pk16_after_xdl
+ ; GFX90A: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8, $vgpr9
+ ; GFX90A-NEXT: {{ $}}
+ ; GFX90A-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr10 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr11 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr12 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr13 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: S_ENDPGM 0
+ $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr10 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr11 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr12 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr13 = nofpexcept V_PK_ADD_F16 8, $vgpr8, 8, $vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ S_ENDPGM 0
+...
+
+---
+name: test_no_unpack_dot_after_xdl
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8, $vgpr9, $vgpr14
+ ; GFX950-LABEL: name: test_no_unpack_dot_after_xdl
+ ; GFX950: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8, $vgpr9, $vgpr14
+ ; GFX950-NEXT: {{ $}}
+ ; GFX950-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr10 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr11 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr12 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr13 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: S_ENDPGM 0
+ ;
+ ; GFX942-LABEL: name: test_no_unpack_dot_after_xdl
+ ; GFX942: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8, $vgpr9, $vgpr14
+ ; GFX942-NEXT: {{ $}}
+ ; GFX942-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr10 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr11 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr12 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr13 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: S_ENDPGM 0
+ ;
+ ; GFX90A-LABEL: name: test_no_unpack_dot_after_xdl
+ ; GFX90A: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr8, $vgpr9, $vgpr14
+ ; GFX90A-NEXT: {{ $}}
+ ; GFX90A-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr10 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr11 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr12 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: $vgpr13 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX90A-NEXT: S_ENDPGM 0
+ $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr10 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr11 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr12 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ $vgpr13 = nofpexcept V_DOT2_F32_F16 8, $vgpr8, 8, $vgpr9, 8, $vgpr14, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ S_ENDPGM 0
+...
>From 5933406dbbf210314035dc43983a23bfb5514c70 Mon Sep 17 00:00:00 2001
From: Li Jinpei <jinpli at amd.com>
Date: Mon, 19 Oct 2026 22:37:20 +0800
Subject: [PATCH 3/3] [AMDGPU] Fix the packed F32 unpack window in
SIPreEmitPeephole
collectUnpackingCandidates unpacked a v_pk_*_f32 only if its two halves
finished before NumPasses - 1, i.e. before the MFMA completes. A packed F32
cannot issue until NumPasses + 2 after an XDL (NumPasses + 3 after an XDL2x),
though, so a packed instruction placed between those points still stalled.
Unpacking pays off whenever the pair finishes before that cycle, so use
getPackedVALUIssueCycle - getVALUIssueCycle + 1 as the window.
Also:
- Skip SGEMM and DGEMM. They hold every VALU until NumPasses + 2, so
unpacking after them only adds an issue slot.
- Charge S_NOP its wait states when counting cycles. The larger window
would otherwise unpack instructions that issue long after the MFMA.
Timed with s_memtime on MI308X (peak TOPS limiter enabled) and MI350X, in
clocks per MFMA followed by four v_pk_fma_f32 (six for F8F6F4):
- XDL 16x16x16: 36 -> 32
- BF16 XDL2x 16x16x32: 40 -> 36
- F8F6F4 XDL2x 16x16x128: 56 -> 52
- SGEMM 16x16x4: 68 -> 56
- DGEMM 16x16x4: 68 -> 56 (MI308X), 104 -> 88 (MI350X)
Packed 16-bit instructions and DOTs are never unpacked, so they don't
change: after an XDL 16x16x16, four v_pk_add_f16 take 20 clocks per MFMA
and four v_dot2_f32_f16 take 40.
With the limiter disabled, MI308X never co-issues VALU with an XDL. The XDL
case becomes 36 -> 40, and four v_pk_add_f16 take 32.
Assisted-by: Claude Opus 5.5
---
llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp | 33 ++++++------
.../AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir | 3 +-
.../AMDGPU/rewrite-mfma-single-exit-copy.ll | 12 +++--
...coissue-insts-post-ra-scheduler-gfx950.mir | 6 ++-
...ck-non-coissue-insts-post-ra-scheduler.mir | 51 ++++++++-----------
5 files changed, 53 insertions(+), 52 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp b/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
index 8825b073ebc40..009ed0949910e 100644
--- a/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp
@@ -19,6 +19,7 @@
//===----------------------------------------------------------------------===//
#include "AMDGPU.h"
+#include "AMDGPUCoExecInfo.h"
#include "GCNSubtarget.h"
#include "llvm/ADT/Statistic.h"
#include "llvm/CodeGen/MachineDominators.h"
@@ -65,10 +66,10 @@ class SIPreEmitPeephole {
bool removeExeczBranch(MachineInstr &MI, MachineBasicBlock &SrcMBB);
bool removeRedundantModeWrites(MachineBasicBlock &SrcMBB) const;
// Creates a list of packed instructions following an MFMA that are suitable
- // for unpacking.
+ // for unpacking, i.e. that would still stall within \p UnpackWindow cycles.
void collectUnpackingCandidates(MachineInstr &BeginMI,
SetVector<MachineInstr *> &InstrsToUnpack,
- uint16_t NumMFMACycles);
+ int UnpackWindow);
// v_pk_fma_f32 v[0:1], v[0:1], v[2:3], v[2:3] op_sel:[1,1,1]
// op_sel_hi:[0,0,0]
// ==>
@@ -711,7 +712,7 @@ void SIPreEmitPeephole::addOperandAndMods(MachineInstrBuilder &NewMI,
void SIPreEmitPeephole::collectUnpackingCandidates(
MachineInstr &BeginMI, SetVector<MachineInstr *> &InstrsToUnpack,
- uint16_t NumMFMACycles) {
+ int UnpackWindow) {
auto *BB = BeginMI.getParent();
auto E = BB->end();
int TotalCyclesBetweenCandidates = 0;
@@ -733,11 +734,13 @@ void SIPreEmitPeephole::collectUnpackingCandidates(
const MCSchedClassDesc *InstrSchedClassDesc =
SchedModel.resolveSchedClass(&Instr);
- uint16_t Latency =
- SchedModel.getWriteProcResBegin(InstrSchedClassDesc)->ReleaseAtCycle;
+ // S_NOP occupies one cycle per wait state.
+ unsigned Latency = std::max<unsigned>(
+ SchedModel.getWriteProcResBegin(InstrSchedClassDesc)->ReleaseAtCycle,
+ SIInstrInfo::getNumWaitStates(Instr));
TotalCyclesBetweenCandidates += Latency;
- if (TotalCyclesBetweenCandidates >= NumMFMACycles - 1)
+ if (TotalCyclesBetweenCandidates >= UnpackWindow)
return;
// Identify register dependencies between those used by the MFMA
// instruction and the following packed instructions. Also checks for
@@ -761,8 +764,7 @@ void SIPreEmitPeephole::collectUnpackingCandidates(
TotalCyclesBetweenCandidates -= Latency;
// TODO: improve latency handling based on instruction modeling.
TotalCyclesBetweenCandidates += 2;
- // Subtract 1 to account for MFMA issue latency.
- if (TotalCyclesBetweenCandidates < NumMFMACycles - 1)
+ if (TotalCyclesBetweenCandidates < UnpackWindow)
InstrsToUnpack.insert(&Instr);
}
}
@@ -914,16 +916,17 @@ bool SIPreEmitPeephole::run(MachineFunction &MF, MachineLoopInfo *LoopInfo) {
for (MachineBasicBlock &MBB : MF) {
// Unpack packed instructions overlapped by MFMAs. This allows the
// compiler to co-issue unpacked instructions with MFMA
- const TargetSchedModel &SchedModel = TII->getSchedModel();
SetVector<MachineInstr *> InstrsToUnpack;
for (auto &MI : make_early_inc_range(MBB.instrs())) {
- if (!SIInstrInfo::isMFMA(MI))
+ // SGEMM/DGEMM hold every VALU back, so only XDL benefits.
+ if (!SIInstrInfo::isMFMA(MI) || !TII->isXDL(MI))
continue;
- const MCSchedClassDesc *SchedClassDesc =
- SchedModel.resolveSchedClass(&MI);
- uint16_t NumMFMACycles =
- SchedModel.getWriteProcResBegin(SchedClassDesc)->ReleaseAtCycle;
- collectUnpackingCandidates(MI, InstrsToUnpack, NumMFMACycles);
+ unsigned PackedCycle =
+ AMDGPU::getPackedVALUIssueCycle(MI, AMDGPU::V_PK_FMA_F32, *TII);
+ unsigned VALUCycle = AMDGPU::getVALUIssueCycle(MI, *TII);
+ assert(PackedCycle >= VALUCycle && "packed VALU issues before VALU");
+ collectUnpackingCandidates(MI, InstrsToUnpack,
+ PackedCycle - VALUCycle + 1);
}
for (MachineInstr *MI : InstrsToUnpack) {
performF32Unpacking(*MI);
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir
index 20a51fd0cc575..6c2eaa128f1ce 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir
@@ -582,7 +582,8 @@
; GCN-NEXT: v_mul_f32_e64 v31, v31, v112
; GCN-NEXT: v_mul_f32_e64 v48, v48, v112
; GCN-NEXT: v_mul_f32_e64 v49, v49, v112
- ; GCN-NEXT: v_pk_mul_f32 v[50:51], v[50:51], v[112:113] op_sel_hi:[1,0]
+ ; GCN-NEXT: v_mul_f32_e64 v50, v50, v112
+ ; GCN-NEXT: v_mul_f32_e64 v51, v51, v112
; GCN-NEXT: v_pk_mul_f32 v[52:53], v[52:53], v[112:113] op_sel_hi:[1,0]
; GCN-NEXT: v_pk_mul_f32 v[54:55], v[54:55], v[112:113] op_sel_hi:[1,0]
; GCN-NEXT: v_pk_mul_f32 v[56:57], v[56:57], v[112:113] op_sel_hi:[1,0]
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-single-exit-copy.ll b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-single-exit-copy.ll
index f0cf067321142..57f1afb367870 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-single-exit-copy.ll
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-single-exit-copy.ll
@@ -304,17 +304,21 @@ define amdgpu_kernel void @single_exit_copy(
; CHECK-NEXT: s_nop 1
; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], a[4:7], a[8:11], a[0:3]
; CHECK-NEXT: s_add_i32 s1, s1, 1
-; CHECK-NEXT: v_pk_add_f32 v[242:243], v[242:243], v[242:243]
-; CHECK-NEXT: v_pk_add_f32 v[240:241], v[240:241], v[240:241]
+; CHECK-NEXT: v_add_f32_e64 v242, v242, v242
+; CHECK-NEXT: v_add_f32_e64 v243, v243, v243
+; CHECK-NEXT: v_add_f32_e64 v240, v240, v240
+; CHECK-NEXT: v_add_f32_e64 v241, v241, v241
; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], a[4:7], a[8:11], a[0:3]
; CHECK-NEXT: v_add_f32_e64 v246, v246, v246
; CHECK-NEXT: v_add_f32_e64 v247, v247, v247
-; CHECK-NEXT: v_pk_add_f32 v[244:245], v[244:245], v[244:245]
+; CHECK-NEXT: v_add_f32_e64 v244, v244, v244
+; CHECK-NEXT: v_add_f32_e64 v245, v245, v245
; CHECK-NEXT: v_pk_add_f32 v[250:251], v[250:251], v[250:251]
; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], a[4:7], a[8:11], a[0:3]
; CHECK-NEXT: v_add_f32_e64 v248, v248, v248
; CHECK-NEXT: v_add_f32_e64 v249, v249, v249
-; CHECK-NEXT: v_pk_add_f32 v[254:255], v[254:255], v[254:255]
+; CHECK-NEXT: v_add_f32_e64 v254, v254, v254
+; CHECK-NEXT: v_add_f32_e64 v255, v255, v255
; CHECK-NEXT: v_pk_add_f32 v[252:253], v[252:253], v[252:253]
; CHECK-NEXT: v_pk_add_f32 v[238:239], v[238:239], v[238:239]
; CHECK-NEXT: v_pk_add_f32 v[236:237], v[236:237], v[236:237]
diff --git a/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler-gfx950.mir b/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler-gfx950.mir
index 96313d997526b..dca70b4efe5b6 100644
--- a/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler-gfx950.mir
+++ b/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler-gfx950.mir
@@ -13,7 +13,8 @@ body: |
; GFX950-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X32_BF16_vgprcd_e64 $vgpr20_vgpr21_vgpr22_vgpr23, $vgpr24_vgpr25_vgpr26_vgpr27, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr12 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr13 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: S_ENDPGM 0
@@ -40,7 +41,8 @@ body: |
; GFX950-NEXT: $vgpr13 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: $vgpr14 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: $vgpr15 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr16 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr17 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: $vgpr18_vgpr19 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: $vgpr20_vgpr21 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: S_ENDPGM 0
diff --git a/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler.mir b/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler.mir
index b56933d627329..0d2fef130de2e 100644
--- a/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler.mir
+++ b/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler.mir
@@ -328,7 +328,8 @@ body: |
; GFX950-NEXT: renamable $vgpr7 = nofpexcept V_ADD_F32_e64 0, killed $sgpr7, 0, $vgpr1, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: renamable $vgpr4 = nofpexcept V_ADD_F32_e64 0, killed $sgpr4, 0, $vgpr0, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: renamable $vgpr5 = nofpexcept V_ADD_F32_e64 0, killed $sgpr5, 0, $vgpr1, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: renamable $vgpr10_vgpr11 = nofpexcept V_PK_ADD_F32 8, killed $sgpr10_sgpr11, 8, $vgpr0_vgpr1, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: renamable $vgpr10 = nofpexcept V_ADD_F32_e64 0, killed $sgpr10, 0, $vgpr0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: renamable $vgpr11 = nofpexcept V_ADD_F32_e64 0, killed $sgpr11, 0, $vgpr1, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: S_ENDPGM 0
;
; GFX942-LABEL: name: test_pk_add_unpacking_f32
@@ -367,7 +368,8 @@ body: |
; GFX942-NEXT: renamable $vgpr7 = nofpexcept V_ADD_F32_e64 0, killed $sgpr7, 0, $vgpr1, 0, 0, implicit $mode, implicit $exec
; GFX942-NEXT: renamable $vgpr4 = nofpexcept V_ADD_F32_e64 0, killed $sgpr4, 0, $vgpr0, 0, 0, implicit $mode, implicit $exec
; GFX942-NEXT: renamable $vgpr5 = nofpexcept V_ADD_F32_e64 0, killed $sgpr5, 0, $vgpr1, 0, 0, implicit $mode, implicit $exec
- ; GFX942-NEXT: renamable $vgpr10_vgpr11 = nofpexcept V_PK_ADD_F32 8, killed $sgpr10_sgpr11, 8, $vgpr0_vgpr1, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: renamable $vgpr10 = nofpexcept V_ADD_F32_e64 0, killed $sgpr10, 0, $vgpr0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: renamable $vgpr11 = nofpexcept V_ADD_F32_e64 0, killed $sgpr11, 0, $vgpr1, 0, 0, implicit $mode, implicit $exec
; GFX942-NEXT: S_ENDPGM 0
;
; GFX90A-LABEL: name: test_pk_add_unpacking_f32
@@ -1180,7 +1182,8 @@ body: |
; GFX950-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr12 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr13 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: S_ENDPGM 0
@@ -1191,7 +1194,8 @@ body: |
; GFX942-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X16BF16_1K_vgprcd_e64 $vgpr4_vgpr5, $vgpr6_vgpr7, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
; GFX942-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
; GFX942-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
- ; GFX942-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr12 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr13 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
; GFX942-NEXT: $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX942-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX942-NEXT: S_ENDPGM 0
@@ -1264,26 +1268,19 @@ body: |
; GFX950: liveins: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7, $vgpr30_vgpr31, $vgpr8_vgpr9
; GFX950-NEXT: {{ $}}
; GFX950-NEXT: early-clobber $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7 = V_MFMA_F64_16X16X4F64_vgprcd_e64 $vgpr30_vgpr31, $vgpr30_vgpr31, $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7, 0, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr12 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr13 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr14 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr15 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr16 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr17 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: S_ENDPGM 0
;
; GFX942-LABEL: name: test_no_unpack_after_dgemm
; GFX942: liveins: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7, $vgpr30_vgpr31, $vgpr8_vgpr9
; GFX942-NEXT: {{ $}}
; GFX942-NEXT: early-clobber $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7 = V_MFMA_F64_16X16X4F64_vgprcd_e64 $vgpr30_vgpr31, $vgpr30_vgpr31, $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7, 0, 0, 0, implicit $mode, implicit $exec
- ; GFX942-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
- ; GFX942-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
- ; GFX942-NEXT: $vgpr12 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
- ; GFX942-NEXT: $vgpr13 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
- ; GFX942-NEXT: $vgpr14 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
- ; GFX942-NEXT: $vgpr15 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX942-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX942-NEXT: S_ENDPGM 0
;
@@ -1314,12 +1311,9 @@ body: |
; GFX950: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr30, $vgpr31, $vgpr8_vgpr9
; GFX950-NEXT: {{ $}}
; GFX950-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X4F32_vgprcd_e64 $vgpr30, $vgpr31, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr12 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr13 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr14 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
- ; GFX950-NEXT: $vgpr15 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX950-NEXT: $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX950-NEXT: S_ENDPGM 0
;
@@ -1327,12 +1321,9 @@ body: |
; GFX942: liveins: $vgpr0_vgpr1_vgpr2_vgpr3, $vgpr30, $vgpr31, $vgpr8_vgpr9
; GFX942-NEXT: {{ $}}
; GFX942-NEXT: $vgpr0_vgpr1_vgpr2_vgpr3 = V_MFMA_F32_16X16X4F32_vgprcd_e64 $vgpr30, $vgpr31, $vgpr0_vgpr1_vgpr2_vgpr3, 0, 0, 0, implicit $mode, implicit $exec
- ; GFX942-NEXT: $vgpr10 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
- ; GFX942-NEXT: $vgpr11 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
- ; GFX942-NEXT: $vgpr12 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
- ; GFX942-NEXT: $vgpr13 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
- ; GFX942-NEXT: $vgpr14 = nofpexcept V_FMA_F32_e64 0, $vgpr8, 0, $vgpr8, 0, $vgpr8, 0, 0, implicit $mode, implicit $exec
- ; GFX942-NEXT: $vgpr15 = nofpexcept V_FMA_F32_e64 0, $vgpr9, 0, $vgpr9, 0, $vgpr9, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr10_vgpr11 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr12_vgpr13 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; GFX942-NEXT: $vgpr14_vgpr15 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX942-NEXT: $vgpr16_vgpr17 = nofpexcept V_PK_FMA_F32 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 8, $vgpr8_vgpr9, 0, 0, 0, 0, 0, implicit $mode, implicit $exec
; GFX942-NEXT: S_ENDPGM 0
;
More information about the llvm-branch-commits
mailing list