[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
Tue Oct 20 00:00:34 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: lijinpei-amd
<details>
<summary>Changes</summary>
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
---
Patch is 49.36 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/233721.diff
8 Files Affected:
- (modified) llvm/lib/Target/AMDGPU/AMDGPUCoExecInfo.h (+12)
- (modified) llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp (+52)
- (modified) llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp (+18-15)
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir (+2-1)
- (modified) llvm/test/CodeGen/AMDGPU/rewrite-mfma-single-exit-copy.ll (+8-4)
- (added) llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler-gfx950.mir (+57)
- (modified) llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler.mir (+264-2)
- (modified) llvm/unittests/Target/AMDGPU/MFMACoExecRules.cpp (+162-2)
``````````diff
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/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
new file mode 100644
index 0000000000000..dca70b4efe5b6
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/unpack-non-coissue-insts-post-ra-scheduler-gfx950.mir
@@ -0,0 +1,57 @@
+# 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 = 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
+ $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 = 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
+ $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..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
@@ -1167,3 +1169,263 @@ 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 = 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
+ ;...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/233721
More information about the llvm-branch-commits
mailing list