[llvm] [AMDGPU] Fix RewriteMFMAFormSchedStage (PR #194887)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Apr 29 18:08:12 PDT 2026
https://github.com/xgxanq updated https://github.com/llvm/llvm-project/pull/194887
>From 2f49c853308448fec1194bb0c7fd1fc18ba4b956 Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Wed, 29 Apr 2026 15:10:42 +0000
Subject: [PATCH 1/2] [AMDGPU] Fix RewriteMFMAFormSchedStage
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 117 +++---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.h | 8 +-
.../rewrite-mfma-form-check-half-rewrite.mir | 86 ++++
.../rewrite-mfma-form-v7slice-pattern.ll | 380 ++++++++++++++++++
4 files changed, 542 insertions(+), 49 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-check-half-rewrite.mir
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-v7slice-pattern.ll
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index af2b2188c3081..d801bceaee09b 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -2250,6 +2250,30 @@ void GCNSchedStage::modifyRegionSchedule(unsigned RegionIdx,
DAG.Regions[RegionIdx].first = MIOrder.front();
}
+void RewriteMFMAFormStage::resetRewriteCandsToVGPR(
+ ArrayRef<std::pair<MachineInstr *, unsigned>> RewriteCands) {
+ for (auto &[MI, OriginalOpcode] : RewriteCands) {
+ assert(TII->isMAI(*MI));
+ const TargetRegisterClass *VDefRC =
+ TII->getRegClass(TII->get(OriginalOpcode), 0);
+ DAG.MRI.setRegClass(MI->getOperand(0).getReg(), VDefRC);
+ MI->setDesc(TII->get(OriginalOpcode));
+
+ MachineOperand *Src2 = TII->getNamedOperand(*MI, AMDGPU::OpName::src2);
+ if (!Src2->isReg())
+ continue;
+
+ // Have to get src types separately since subregs may cause C and D
+ // registers to be different types even though the actual operand is
+ // the same size.
+ const TargetRegisterClass *AUseRC =
+ DAG.MRI.getRegClass(Src2->getReg());
+ const TargetRegisterClass *VUseRC =
+ SRI->getEquivalentVGPRClass(AUseRC);
+ DAG.MRI.setRegClass(Src2->getReg(), VUseRC);
+ }
+}
+
bool RewriteMFMAFormStage::isRewriteCandidate(MachineInstr *MI) const {
if (!static_cast<const SIInstrInfo *>(DAG.TII)->isMAI(*MI))
@@ -2272,10 +2296,18 @@ bool RewriteMFMAFormStage::initHeuristics(
int ReplacementOp = AMDGPU::getMFMASrcCVDstAGPROp(MI.getOpcode());
assert(ReplacementOp != -1);
+ MachineOperand *Src2 = TII->getNamedOperand(MI, AMDGPU::OpName::src2);
+ MachineOperand &Dst = MI.getOperand(0);
+ assert(Src2);
+ // Pre-validate: both dst and src2 (if a register) must be virtual.
+ if (!Dst.getReg().isVirtual() ||
+ (Src2->isReg() && !Src2->getReg().isVirtual())) {
+ continue;
+ }
+
RewriteCands.push_back({&MI, MI.getOpcode()});
MI.setDesc(TII->get(ReplacementOp));
- MachineOperand *Src2 = TII->getNamedOperand(MI, AMDGPU::OpName::src2);
if (Src2->isReg()) {
SmallVector<SlotIndex, 8> Src2ReachingDefs;
findReachingDefs(*Src2, DAG.LIS, Src2ReachingDefs);
@@ -2289,7 +2321,6 @@ bool RewriteMFMAFormStage::initHeuristics(
}
}
- MachineOperand &Dst = MI.getOperand(0);
SmallVector<MachineOperand *, 8> DstReachingUses;
findReachingUses(&MI, DAG.LIS, DstReachingUses);
@@ -2339,7 +2370,7 @@ bool RewriteMFMAFormStage::initHeuristics(
}
int64_t RewriteMFMAFormStage::getRewriteCost(
- const std::vector<std::pair<MachineInstr *, unsigned>> &RewriteCands,
+ ArrayRef<std::pair<MachineInstr *, unsigned>> RewriteCands,
const DenseMap<MachineBasicBlock *, std::set<Register>> &CopyForUse,
const SmallPtrSetImpl<MachineInstr *> &CopyForDef) {
MachineBlockFrequencyInfo *MBFI = DAG.MBFI;
@@ -2354,6 +2385,10 @@ int64_t RewriteMFMAFormStage::getRewriteCost(
unsigned AGPRThreshold = MaxVectorRegs.second;
unsigned CombinedThreshold = ST.getMaxNumVGPRs(MF);
+ // Reset the classes that were changed to AGPR for better RB analysis.
+ // We must do rewriting after copy-insertion, as some defs of the register
+ // may require VGPR. Additionally, if we bail out and don't perform the
+ // rewrite then these need to be restored anyway.
for (unsigned Region = 0; Region < DAG.Regions.size(); Region++) {
if (!RegionsWithExcessArchVGPR[Region])
continue;
@@ -2391,8 +2426,10 @@ int64_t RewriteMFMAFormStage::getRewriteCost(
SpillCost *= (int64_t)RelativeFreq;
// If we have increased spilling in any block, just bail.
- if (SpillCost > 0)
+ if (SpillCost > 0) {
+ resetRewriteCandsToVGPR(RewriteCands);
return SpillCost;
+ }
if (SpillCost < BestSpillCost)
BestSpillCost = SpillCost;
@@ -2429,36 +2466,13 @@ int64_t RewriteMFMAFormStage::getRewriteCost(
}
}
- // Reset the classes that were changed to AGPR for better RB analysis.
- // We must do rewriting after copy-insertion, as some defs of the register
- // may require VGPR. Additionally, if we bail out and don't perform the
- // rewrite then these need to be restored anyway.
- for (auto &[MI, OriginalOpcode] : RewriteCands) {
- assert(TII->isMAI(*MI));
- const TargetRegisterClass *ADefRC =
- DAG.MRI.getRegClass(MI->getOperand(0).getReg());
- const TargetRegisterClass *VDefRC = SRI->getEquivalentVGPRClass(ADefRC);
- DAG.MRI.setRegClass(MI->getOperand(0).getReg(), VDefRC);
- MI->setDesc(TII->get(OriginalOpcode));
-
- MachineOperand *Src2 = TII->getNamedOperand(*MI, AMDGPU::OpName::src2);
- assert(Src2);
- if (!Src2->isReg())
- continue;
-
- // Have to get src types separately since subregs may cause C and D
- // registers to be different types even though the actual operand is
- // the same size.
- const TargetRegisterClass *AUseRC = DAG.MRI.getRegClass(Src2->getReg());
- const TargetRegisterClass *VUseRC = SRI->getEquivalentVGPRClass(AUseRC);
- DAG.MRI.setRegClass(Src2->getReg(), VUseRC);
- }
+ resetRewriteCandsToVGPR(RewriteCands);
return Cost + CopyCost;
}
bool RewriteMFMAFormStage::rewrite(
- const std::vector<std::pair<MachineInstr *, unsigned>> &RewriteCands) {
+ ArrayRef<std::pair<MachineInstr *, unsigned>> RewriteCands) {
DenseMap<MachineInstr *, unsigned> FirstMIToRegion;
DenseMap<MachineInstr *, unsigned> LastMIToRegion;
@@ -2534,9 +2548,6 @@ bool RewriteMFMAFormStage::rewrite(
MachineOperand *Src2 = TII->getNamedOperand(*MI, AMDGPU::OpName::src2);
if (Src2->isReg()) {
Register Src2Reg = Src2->getReg();
- if (!Src2Reg.isVirtual())
- return false;
-
Register MappedReg = Src2->getReg();
SmallVector<SlotIndex, 8> Src2ReachingDefs;
findReachingDefs(*Src2, DAG.LIS, Src2ReachingDefs);
@@ -2602,9 +2613,6 @@ bool RewriteMFMAFormStage::rewrite(
MachineOperand *Dst = &MI->getOperand(0);
Register DstReg = Dst->getReg();
- if (!DstReg.isVirtual())
- return false;
-
Register MappedReg = DstReg;
SmallVector<MachineOperand *, 8> DstReachingUses;
@@ -2675,6 +2683,14 @@ bool RewriteMFMAFormStage::rewrite(
}
DenseSet<MachineOperand *> &DstRegSet = ReplaceMap[DstReg];
+
+ // For same-block sub-register uses: all RUs that share the same DstReg
+ // and are in the same MBB as the MFMA can reuse a single COPY instead of
+ // inserting one COPY per use. The COPY is placed immediately after the
+ // MFMA so it dominates every same-block use regardless of their order in
+ // DstReachingUseCopies.
+ Register SameBlockNewUseReg;
+
for (MachineOperand *RU : DstReachingUseCopies) {
MachineBasicBlock *RUBlock = RU->getParent()->getParent();
// Just keep track of the reaching use of this register by block. After we
@@ -2684,22 +2700,29 @@ bool RewriteMFMAFormStage::rewrite(
continue;
}
- // Special case, the use is in the same block as the MFMA. Insert the copy
- // just before the use.
+ // Special case: the use is in the same block as the MFMA. Insert a single
+ // COPY immediately after the MFMA and reuse it for all same-block uses.
+ if (SameBlockNewUseReg.isValid()) {
+ // Reuse the COPY already inserted for this DstReg; no new instruction.
+ RU->setReg(SameBlockNewUseReg);
+ continue;
+ }
+
+ // First same-block use: create one COPY placed right after the MFMA so
+ // it dominates all subsequent same-block uses of DstReg.
const TargetRegisterClass *DstRC = DAG.MRI.getRegClass(DstReg);
const TargetRegisterClass *VGPRRC = SRI->getEquivalentVGPRClass(DstRC);
- Register NewUseReg = DAG.MRI.createVirtualRegister(VGPRRC);
- MachineInstr *UseInst = RU->getParent();
+ SameBlockNewUseReg = DAG.MRI.createVirtualRegister(VGPRRC);
MachineInstrBuilder VGPRCopy =
- BuildMI(*UseInst->getParent(), UseInst->getIterator(),
- UseInst->getDebugLoc(), TII->get(TargetOpcode::COPY))
- .addDef(NewUseReg, {}, 0)
+ BuildMI(*MI->getParent(), std::next(MI->getIterator()),
+ MI->getDebugLoc(), TII->get(TargetOpcode::COPY))
+ .addDef(SameBlockNewUseReg, {}, 0)
.addUse(DstReg, {}, 0);
DAG.LIS->InsertMachineInstrInMaps(*VGPRCopy);
// Since we know this use has only one reaching def, we can replace the
// use reg.
- RU->setReg(NewUseReg);
- // Track the copy source operand for r eplacement.
+ RU->setReg(SameBlockNewUseReg);
+ // Track the copy source operand for replacement.
DstRegSet.insert(&VGPRCopy->getOperand(1));
}
@@ -2793,10 +2816,10 @@ bool RewriteMFMAFormStage::rewrite(
RegionPressureMap LiveInUpdater(&DAG, false);
LiveInUpdater.buildLiveRegMap();
- for (unsigned Region = 0; Region < DAG.Regions.size(); Region++)
+ for (unsigned Region = 0; Region < DAG.Regions.size(); Region++) {
DAG.LiveIns[Region] = LiveInUpdater.getLiveRegsForRegionIdx(Region);
-
- DAG.Pressure[RegionIdx] = DAG.getRealRegPressure(RegionIdx);
+ DAG.Pressure[Region] = DAG.getRealRegPressure(Region);
+ }
return true;
}
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
index 2cc9e81a65191..684267d381478 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
@@ -459,13 +459,17 @@ class RewriteMFMAFormStage : public GCNSchedStage {
/// in initHeuristics. Uses \p CopyForUse and \p CopyForDef to calculate copy
/// costs, and \p RewriteCands to undo rewriting.
int64_t getRewriteCost(
- const std::vector<std::pair<MachineInstr *, unsigned>> &RewriteCands,
+ ArrayRef<std::pair<MachineInstr *, unsigned>> RewriteCands,
const DenseMap<MachineBasicBlock *, std::set<Register>> &CopyForUse,
const SmallPtrSetImpl<MachineInstr *> &CopyForDef);
/// Do the final rewrite on \p RewriteCands and insert any needed copies.
bool
- rewrite(const std::vector<std::pair<MachineInstr *, unsigned>> &RewriteCands);
+ rewrite(ArrayRef<std::pair<MachineInstr *, unsigned>> RewriteCands);
+ /// Resets all rewrite candidates in \p Cands back to their original VGPR
+ /// opcodes and register classes.
+ void resetRewriteCandsToVGPR(
+ ArrayRef<std::pair<MachineInstr *, unsigned>> RewriteCands);
/// \returns true if this MI is a rewrite candidate.
bool isRewriteCandidate(MachineInstr *MI) const;
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-check-half-rewrite.mir b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-check-half-rewrite.mir
new file mode 100644
index 0000000000000..3a73e82395b9b
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-check-half-rewrite.mir
@@ -0,0 +1,86 @@
+# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 \
+# RUN: -run-pass=machine-scheduler \
+# RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+# RUN: -o - %s | FileCheck %s
+#
+# Test: RewriteMFMAFormStage — physical register in MFMA src2 aborts rewrite cleanly.
+#
+# Root cause (GCNSchedStrategy.cpp):
+# rewrite() processes RewriteCands in iteration order. For each MFMA:
+# MI->setDesc(ReplacementOp); // opcode mutated
+# if (!Src2Reg.isVirtual()) return false; // mid-loop bail-out on physreg
+# Fix: check isVirtual() in isRewriteCandidate() before adding to RewriteCands,
+# so rewrite() never sees a physreg src2. Both MFMAs stay in VGPR form.
+#
+# Trigger:
+# MFMA_B uses $vgpr4_vgpr5_vgpr6_vgpr7 (physical VGPR) as src2.
+# With fix: isRewriteCandidate() rejects MFMA_B -> rewrite() not called ->
+# no crash, both MFMAs remain in vreg_128_align2 / vgprcd form.
+#
+# CHECK-LABEL: name: test_bug2_physreg_src2
+# CHECK: bb.0:
+# MFMA_A: virtual src2 but rewrite aborted (MFMA_B not a candidate) -> vgprcd stays.
+# CHECK: %res_a:vreg_128_align2 = {{.*}} V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64
+# MFMA_B: physical src2 — skipped by isRewriteCandidate, stays in vgprcd form.
+# CHECK: %res_b:vreg_128_align2 = {{.*}} V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64
+
+--- |
+ define void @test_bug2_physreg_src2() #0 {
+ entry:
+ unreachable
+ }
+
+ attributes #0 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" }
+...
+
+---
+name: test_bug2_physreg_src2
+tracksRegLiveness: true
+liveins:
+ - { reg: '$vgpr4' }
+ - { reg: '$vgpr5' }
+ - { reg: '$vgpr6' }
+ - { reg: '$vgpr7' }
+machineFunctionInfo:
+ isEntryFunction: true
+ scratchRSrcReg: '$sgpr96_sgpr97_sgpr98_sgpr99'
+ stackPtrOffsetReg: '$sgpr32'
+ argumentInfo:
+ privateSegmentBuffer: { reg: '$sgpr0_sgpr1_sgpr2_sgpr3' }
+ kernargSegmentPtr: { reg: '$sgpr4_sgpr5' }
+ workGroupIDX: { reg: '$sgpr6' }
+ privateSegmentWaveByteOffset: { reg: '$sgpr7' }
+ workItemIDX: { reg: '$vgpr0' }
+ sgprForEXECCopy: '$sgpr100_sgpr101'
+body: |
+ bb.0:
+ liveins: $vgpr0, $sgpr4_sgpr5, $vgpr4, $vgpr5, $vgpr6, $vgpr7
+
+ ; 9 x vreg_1024 = 288 VGPRs > 256: triggers RegionsWithExcessArchVGPR.
+ %0:vreg_1024 = IMPLICIT_DEF
+ %1:vreg_1024 = IMPLICIT_DEF
+ %2:vreg_1024 = IMPLICIT_DEF
+ %3:vreg_1024 = IMPLICIT_DEF
+ %4:vreg_1024 = IMPLICIT_DEF
+ %5:vreg_1024 = IMPLICIT_DEF
+ %6:vreg_1024 = IMPLICIT_DEF
+ %7:vreg_1024 = IMPLICIT_DEF
+ %8:vreg_1024 = IMPLICIT_DEF
+
+ %9:av_128_align2 = IMPLICIT_DEF
+ %10:av_128_align2 = IMPLICIT_DEF
+ %11:vreg_64_align2 = IMPLICIT_DEF
+ %12:vgpr_32 = IMPLICIT_DEF
+
+ SCHED_BARRIER 0
+
+ ; MFMA_A: all virtual registers.
+ %acc_a:vreg_128_align2 = IMPLICIT_DEF
+ %res_a:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 %9:av_128_align2, %10:av_128_align2, %acc_a:vreg_128_align2, 4, 4, %11.sub0:vreg_64_align2, %12:vgpr_32, 0, 0, implicit $mode, implicit $exec
+
+ ; MFMA_B: src2 = $vgpr4_vgpr5_vgpr6_vgpr7 (physical VGPR) — triggers crash.
+ %res_b:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 %9:av_128_align2, %10:av_128_align2, $vgpr4_vgpr5_vgpr6_vgpr7, 4, 4, %11.sub0:vreg_64_align2, %12:vgpr_32, 0, 0, implicit $mode, implicit $exec
+
+ KILL %0, %1, %2, %3, %4, %5, %6, %7, %8, %res_a, %res_b
+
+ S_ENDPGM 0
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-v7slice-pattern.ll b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-v7slice-pattern.ll
new file mode 100644
index 0000000000000..7804a2f61c4c9
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-v7slice-pattern.ll
@@ -0,0 +1,380 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 \
+; RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+; RUN: < %s | FileCheck %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 \
+; RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+; RUN: -stop-before=machine-scheduler \
+; RUN: < %s | FileCheck %s --check-prefix=BEFORE
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 \
+; RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+; RUN: -stop-after=machine-scheduler \
+; RUN: < %s | FileCheck %s --check-prefix=AFTER
+;
+; Test: RewriteMFMAFormStage — v7_slice pattern (gfx950, mfma.f32.16x16x32.f16)
+;
+; Distilled from v7_slice.llir (Triton matmul kernel, gfx950).
+;
+; Key structural features preserved from v7_slice:
+; (1) Loop body (%loop) accumulates via loop-carried SCALAR float phis,
+; NOT <4 x float> phis. Each MFMA acc is built by insertelement from
+; 4 individual float phis then used as src2/dst in MFMAs within the loop.
+; (2) Epilogue (%epilogue) ALSO contains MFMAs: the final K-tile iteration
+; is peeled. Epilogue MFMAs take scalar float phis from %loop as acc,
+; then produce results consumed by fptrunc+store — Case 2.
+; (3) Scalar float phis initialized to 0.0 in entry (non-MAI def) — Case 3.
+;
+; CFG: entry -> loop (back-edge) -> epilogue -> ret
+;
+; Pressure design (gfx950, 256 ArchVGPR limit):
+; 64 loop-carried scalar float phis = 64 VGPRs (non-MFMA pressure carriers).
+; 4 MFMA chains of 2, each acc = 4 float phis = 16 VGPRs in loop.
+; Plus <8 x half> src0/src1 operands reused across chains = 8 VGPRs.
+; Total in loop body: 64 + 16 + 8 = 88 VGPRs (not enough alone).
+; Add 6 more <32 x float> loop-carried vec carriers = 192 VGPRs.
+; Peak ArchVGPR = 192 + 64 + 16 + 8 = 280 > 256 -> RegionsWithExcessArchVGPR.
+; After rewrite: scalar phis=64 VGPR, vec carriers=192 VGPR, MFMA acc->AGPR.
+; getRewriteCost() < 0 -> rewrite() fires.
+;
+; Expected:
+; Case 3: v_accvgpr_write inserted after insertelement defs of MFMA acc
+; (entry zeroinitializer -> scalar float phi -> insertelement = non-MAI)
+; Case 2: v_accvgpr_read before fptrunc in epilogue block
+; MFMA opcodes: v_mfma_f32_16x16x32_f16 vgprcd -> agprcd form
+
+declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half>, <8 x half>, <4 x float>, i32 immarg, i32 immarg, i32 immarg)
+
+define amdgpu_kernel void @test_v7slice_scalar_phi_acc(
+; CHECK-LABEL: test_v7slice_scalar_phi_acc:
+; CHECK: ; @test_v7slice_scalar_phi_acc
+; Case 3: v_accvgpr_write inserted for scalar-phi acc (non-MAI def -> AGPR init).
+; Loop body MFMAs use AGPR C/D (a[...], a[...], a[...], a[...]).
+; CHECK: ; %bb.0: ; %entry
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: .LBB0_1: ; %loop
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; Case 2: epilogue MFMAs take VGPR src0/src1, AGPR C/D. v_accvgpr_read before fptrunc.
+; CHECK: ; %bb.2: ; %epilogue
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+ ptr addrspace(1) %out,
+ <8 x half> %a0,
+ <8 x half> %a1,
+ <8 x half> %b0,
+ <8 x half> %b1,
+ i32 %n) #0 {
+entry:
+ br label %loop
+
+loop:
+ %i = phi i32 [ 0, %entry ], [ %i.next, %loop ]
+
+ ; -------------------------------------------------------------------
+ ; 8 x <32 x float> loop-carried vector carriers = 256 VGPRs.
+ ; Self-insert with variable index prevents folding.
+ ; Kept live by store use in loop body.
+ ; -------------------------------------------------------------------
+ %vc0 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc0n, %loop ]
+ %vc1 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc1n, %loop ]
+ %vc2 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc2n, %loop ]
+ %vc3 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc3n, %loop ]
+ %vc4 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc4n, %loop ]
+ %vc5 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc5n, %loop ]
+ %vc6 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc6n, %loop ]
+ %vc7 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc7n, %loop ]
+
+ ; -------------------------------------------------------------------
+ ; 16 loop-carried scalar float phis = 16 VGPRs.
+ ; These are the per-element accumulator slots, mirroring v7_slice's
+ ; pattern where each output element is a separate loop-carried float.
+ ; Initialized to 0.0 in entry (non-MAI def) -> Case 3 triggers on the
+ ; insertelement that builds the MFMA src2 from these scalars.
+ ; -------------------------------------------------------------------
+ ; MFMA chain A acc (4 floats):
+ %a0_0 = phi float [ 0.0, %entry ], [ %ra0_0, %loop ]
+ %a0_1 = phi float [ 0.0, %entry ], [ %ra0_1, %loop ]
+ %a0_2 = phi float [ 0.0, %entry ], [ %ra0_2, %loop ]
+ %a0_3 = phi float [ 0.0, %entry ], [ %ra0_3, %loop ]
+ ; MFMA chain B acc (4 floats):
+ %a1_0 = phi float [ 0.0, %entry ], [ %ra1_0, %loop ]
+ %a1_1 = phi float [ 0.0, %entry ], [ %ra1_1, %loop ]
+ %a1_2 = phi float [ 0.0, %entry ], [ %ra1_2, %loop ]
+ %a1_3 = phi float [ 0.0, %entry ], [ %ra1_3, %loop ]
+ ; MFMA chain C acc (4 floats):
+ %a2_0 = phi float [ 0.0, %entry ], [ %ra2_0, %loop ]
+ %a2_1 = phi float [ 0.0, %entry ], [ %ra2_1, %loop ]
+ %a2_2 = phi float [ 0.0, %entry ], [ %ra2_2, %loop ]
+ %a2_3 = phi float [ 0.0, %entry ], [ %ra2_3, %loop ]
+ ; MFMA chain D acc (4 floats):
+ %a3_0 = phi float [ 0.0, %entry ], [ %ra3_0, %loop ]
+ %a3_1 = phi float [ 0.0, %entry ], [ %ra3_1, %loop ]
+ %a3_2 = phi float [ 0.0, %entry ], [ %ra3_2, %loop ]
+ %a3_3 = phi float [ 0.0, %entry ], [ %ra3_3, %loop ]
+
+ ; -------------------------------------------------------------------
+ ; Build <4 x float> acc vectors from scalar phi elements (v7_slice pattern).
+ ; Each insertelement is a non-MAI def of an element of the acc register.
+ ; -------------------------------------------------------------------
+ %accA0 = insertelement <4 x float> poison, float %a0_0, i32 0
+ %accA1 = insertelement <4 x float> %accA0, float %a0_1, i32 1
+ %accA2 = insertelement <4 x float> %accA1, float %a0_2, i32 2
+ %accA = insertelement <4 x float> %accA2, float %a0_3, i32 3
+
+ %accB0 = insertelement <4 x float> poison, float %a1_0, i32 0
+ %accB1 = insertelement <4 x float> %accB0, float %a1_1, i32 1
+ %accB2 = insertelement <4 x float> %accB1, float %a1_2, i32 2
+ %accB = insertelement <4 x float> %accB2, float %a1_3, i32 3
+
+ %accC0 = insertelement <4 x float> poison, float %a2_0, i32 0
+ %accC1 = insertelement <4 x float> %accC0, float %a2_1, i32 1
+ %accC2 = insertelement <4 x float> %accC1, float %a2_2, i32 2
+ %accC = insertelement <4 x float> %accC2, float %a2_3, i32 3
+
+ %accD0 = insertelement <4 x float> poison, float %a3_0, i32 0
+ %accD1 = insertelement <4 x float> %accD0, float %a3_1, i32 1
+ %accD2 = insertelement <4 x float> %accD1, float %a3_2, i32 2
+ %accD = insertelement <4 x float> %accD2, float %a3_3, i32 3
+
+ ; -------------------------------------------------------------------
+ ; MFMA chains in loop body (2 MFMAs per chain, chained acc).
+ ; Results are loop-carried back via extractelement -> scalar phi.
+ ; -------------------------------------------------------------------
+ %rA_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %accA, i32 0, i32 0, i32 0)
+ %rA = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %rA_v, i32 0, i32 0, i32 0)
+
+ %rB_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %accB, i32 0, i32 0, i32 0)
+ %rB = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %rB_v, i32 0, i32 0, i32 0)
+
+ %rC_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %accC, i32 0, i32 0, i32 0)
+ %rC = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %rC_v, i32 0, i32 0, i32 0)
+
+ %rD_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %accD, i32 0, i32 0, i32 0)
+ %rD = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %rD_v, i32 0, i32 0, i32 0)
+
+ ; Extract scalar results to carry back via phi.
+ %ra0_0 = extractelement <4 x float> %rA, i32 0
+ %ra0_1 = extractelement <4 x float> %rA, i32 1
+ %ra0_2 = extractelement <4 x float> %rA, i32 2
+ %ra0_3 = extractelement <4 x float> %rA, i32 3
+
+ %ra1_0 = extractelement <4 x float> %rB, i32 0
+ %ra1_1 = extractelement <4 x float> %rB, i32 1
+ %ra1_2 = extractelement <4 x float> %rB, i32 2
+ %ra1_3 = extractelement <4 x float> %rB, i32 3
+
+ %ra2_0 = extractelement <4 x float> %rC, i32 0
+ %ra2_1 = extractelement <4 x float> %rC, i32 1
+ %ra2_2 = extractelement <4 x float> %rC, i32 2
+ %ra2_3 = extractelement <4 x float> %rC, i32 3
+
+ %ra3_0 = extractelement <4 x float> %rD, i32 0
+ %ra3_1 = extractelement <4 x float> %rD, i32 1
+ %ra3_2 = extractelement <4 x float> %rD, i32 2
+ %ra3_3 = extractelement <4 x float> %rD, i32 3
+
+ ; Vector carrier self-inserts (prevent register reuse).
+ %eidx = and i32 %i, 31
+ %vc0e = extractelement <32 x float> %vc0, i32 0
+ %vc1e = extractelement <32 x float> %vc1, i32 0
+ %vc2e = extractelement <32 x float> %vc2, i32 0
+ %vc3e = extractelement <32 x float> %vc3, i32 0
+ %vc4e = extractelement <32 x float> %vc4, i32 0
+ %vc5e = extractelement <32 x float> %vc5, i32 0
+ %vc6e = extractelement <32 x float> %vc6, i32 0
+ %vc7e = extractelement <32 x float> %vc7, i32 0
+ %vc0n = insertelement <32 x float> %vc0, float %vc0e, i32 %eidx
+ %vc1n = insertelement <32 x float> %vc1, float %vc1e, i32 %eidx
+ %vc2n = insertelement <32 x float> %vc2, float %vc2e, i32 %eidx
+ %vc3n = insertelement <32 x float> %vc3, float %vc3e, i32 %eidx
+ %vc4n = insertelement <32 x float> %vc4, float %vc4e, i32 %eidx
+ %vc5n = insertelement <32 x float> %vc5, float %vc5e, i32 %eidx
+ %vc6n = insertelement <32 x float> %vc6, float %vc6e, i32 %eidx
+ %vc7n = insertelement <32 x float> %vc7, float %vc7e, i32 %eidx
+
+ %vcs = fadd float %vc0e, %vc1e
+ store float %vcs, ptr addrspace(1) %out, align 4
+
+ %i.next = add i32 %i, 1
+ %cond = icmp eq i32 %i.next, %n
+ br i1 %cond, label %epilogue, label %loop
+
+epilogue:
+ ; -------------------------------------------------------------------
+ ; v7_slice pattern: epilogue ALSO has MFMAs (peeled final K-tile).
+ ; Acc built from the same scalar float phis coming out of %loop.
+ ; These insertelements are non-MAI defs -> Case 3.
+ ; -------------------------------------------------------------------
+ %eaccA0 = insertelement <4 x float> poison, float %ra0_0, i32 0
+ %eaccA1 = insertelement <4 x float> %eaccA0, float %ra0_1, i32 1
+ %eaccA2 = insertelement <4 x float> %eaccA1, float %ra0_2, i32 2
+ %eaccA = insertelement <4 x float> %eaccA2, float %ra0_3, i32 3
+
+ %eaccB0 = insertelement <4 x float> poison, float %ra1_0, i32 0
+ %eaccB1 = insertelement <4 x float> %eaccB0, float %ra1_1, i32 1
+ %eaccB2 = insertelement <4 x float> %eaccB1, float %ra1_2, i32 2
+ %eaccB = insertelement <4 x float> %eaccB2, float %ra1_3, i32 3
+
+ %eaccC0 = insertelement <4 x float> poison, float %ra2_0, i32 0
+ %eaccC1 = insertelement <4 x float> %eaccC0, float %ra2_1, i32 1
+ %eaccC2 = insertelement <4 x float> %eaccC1, float %ra2_2, i32 2
+ %eaccC = insertelement <4 x float> %eaccC2, float %ra2_3, i32 3
+
+ %eaccD0 = insertelement <4 x float> poison, float %ra3_0, i32 0
+ %eaccD1 = insertelement <4 x float> %eaccD0, float %ra3_1, i32 1
+ %eaccD2 = insertelement <4 x float> %eaccD1, float %ra3_2, i32 2
+ %eaccD = insertelement <4 x float> %eaccD2, float %ra3_3, i32 3
+
+ ; Epilogue MFMAs (2 per chain, as in v7_slice).
+ %erA_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %eaccA, i32 0, i32 0, i32 0)
+ %erA = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %erA_v, i32 0, i32 0, i32 0)
+
+ %erB_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %eaccB, i32 0, i32 0, i32 0)
+ %erB = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %erB_v, i32 0, i32 0, i32 0)
+
+ %erC_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %eaccC, i32 0, i32 0, i32 0)
+ %erC = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %erC_v, i32 0, i32 0, i32 0)
+
+ %erD_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %eaccD, i32 0, i32 0, i32 0)
+ %erD = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %erD_v, i32 0, i32 0, i32 0)
+
+ ; -------------------------------------------------------------------
+ ; Non-MAI use of epilogue MFMA results via fptrunc (Case 2).
+ ; v7_slice: results shuffled and fptrunc'd to fp16 for stores.
+ ; -------------------------------------------------------------------
+ %shA = shufflevector <4 x float> %erA, <4 x float> poison, <2 x i32> <i32 0, i32 1>
+ %hA = fptrunc <2 x float> %shA to <2 x half>
+ store <2 x half> %hA, ptr addrspace(1) %out, align 2
+
+ %shB = shufflevector <4 x float> %erB, <4 x float> poison, <2 x i32> <i32 0, i32 1>
+ %hB = fptrunc <2 x float> %shB to <2 x half>
+ %pB = getelementptr i16, ptr addrspace(1) %out, i32 2
+ store <2 x half> %hB, ptr addrspace(1) %pB, align 2
+
+ %shC = shufflevector <4 x float> %erC, <4 x float> poison, <2 x i32> <i32 0, i32 1>
+ %hC = fptrunc <2 x float> %shC to <2 x half>
+ %pC = getelementptr i16, ptr addrspace(1) %out, i32 4
+ store <2 x half> %hC, ptr addrspace(1) %pC, align 2
+
+ %shD = shufflevector <4 x float> %erD, <4 x float> poison, <2 x i32> <i32 0, i32 1>
+ %hD = fptrunc <2 x float> %shD to <2 x half>
+ %pD = getelementptr i16, ptr addrspace(1) %out, i32 6
+ store <2 x half> %hD, ptr addrspace(1) %pD, align 2
+
+ ret void
+}
+
+attributes #0 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" }
+
+; stop-before=machine-scheduler: MFMAs still in VGPR (vgprcd) form.
+; BEFORE-LABEL: name: test_v7slice_scalar_phi_acc
+; 8x V_MFMA_F32_16X16X32_F16_vgprcd_e64: VGPR form, rewrite not yet applied.
+; BEFORE: bb.1.loop:
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; 8x V_MFMA_F32_16X16X32_F16_vgprcd_e64: epilogue peeled K-tile, VGPR form.
+; BEFORE: bb.2.epilogue:
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; BEFORE: %{{[0-9]+}}:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+
+; AFTER-LABEL: name: test_v7slice_scalar_phi_acc
+; Case 3: 4x areg_128_align2 = COPY vreg (VGPR->AGPR init for scalar-phi acc).
+; AFTER: bb.0.entry:
+; AFTER: %{{[0-9]+}}:areg_128_align2 = COPY %{{[0-9]+}}
+; AFTER: %{{[0-9]+}}:areg_128_align2 = COPY %{{[0-9]+}}
+; AFTER: %{{[0-9]+}}:areg_128_align2 = COPY %{{[0-9]+}}
+; AFTER: %{{[0-9]+}}:areg_128_align2 = COPY %{{[0-9]+}}
+; 8x V_MFMA_F32_16X16X32_F16_e64: AGPR form, loop-carried acc in areg.
+; AFTER: bb.1.loop:
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; Case 2: vreg_128_align2 = COPY areg interleaved with epilogue MFMAs (AGPR->VGPR).
+; AFTER: bb.2.epilogue:
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:vreg_128_align2 = COPY %{{[0-9]+}}
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:vreg_128_align2 = COPY %{{[0-9]+}}
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:vreg_128_align2 = COPY %{{[0-9]+}}
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 %{{[0-9]+}}, %{{[0-9]+}}, %{{[0-9]+}}, 0, 0, 0, implicit $mode, implicit $exec
+; AFTER: %{{[0-9]+}}:vreg_128_align2 = COPY %{{[0-9]+}}
>From 4e5a01ccb54e17e19764df424a4c013133028563 Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Wed, 29 Apr 2026 15:10:42 +0000
Subject: [PATCH 2/2] [AMDGPU] Fix RewriteMFMAFormSchedStage
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 42 +++++++++----------
llvm/lib/Target/AMDGPU/GCNSchedStrategy.h | 24 +++++------
.../rewrite-mfma-form-check-half-rewrite.mir | 7 ++--
3 files changed, 35 insertions(+), 38 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index d801bceaee09b..e376baf118f1d 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -16,7 +16,7 @@
/// GCNScheduleDAGMILive::runSchedStages.
/// Generally, the reason for having multiple scheduling stages is to account
-/// for the kernel-wide effect of register usage on occupancy. Usually, only a
+/// for the kernel-wide effect of register usage on occupancy. Usually, only a
/// few scheduling regions will have register pressure high enough to limit
/// occupancy for the kernel, so constraints can be relaxed to improve ILP in
/// other regions.
@@ -124,7 +124,7 @@ void GCNSchedStrategy::initialize(ScheduleDAGMI *DAG) {
Context->RegClassInfo->getNumAllocatableRegs(&AMDGPU::VGPR_32RegClass);
SIMachineFunctionInfo &MFI = *MF->getInfo<SIMachineFunctionInfo>();
- // Set the initial TargetOccupnacy to the maximum occupancy that we can
+ // Set the initial TargetOccupancy to the maximum occupancy that we can
// achieve for this function. This effectively sets a lower bound on the
// 'Critical' register limits in the scheduler.
// Allow for lower occupancy targets if kernel is wave limited or memory
@@ -339,7 +339,7 @@ void GCNSchedStrategy::initCandidate(SchedCandidate &Cand, SUnit *SU,
// If two instructions increase the pressure of different register sets
// by the same amount, the generic scheduler will prefer to schedule the
// instruction that increases the set with the least amount of registers,
- // which in our case would be SGPRs. This is rarely what we want, so
+ // which in our case would be SGPRs. This is rarely what we want, so
// when we report excess/critical register pressure, we do it either
// only for VGPRs or only for SGPRs.
@@ -349,7 +349,7 @@ void GCNSchedStrategy::initCandidate(SchedCandidate &Cand, SUnit *SU,
bool ShouldTrackSGPRs = !ShouldTrackVGPRs && SGPRPressure >= SGPRExcessLimit;
// FIXME: We have to enter REG-EXCESS before we reach the actual threshold
- // to increase the likelihood we don't go over the limits. We should improve
+ // to increase the likelihood we don't go over the limits. We should improve
// the analysis to look through dependencies to find the path with the least
// register pressure.
@@ -370,7 +370,7 @@ void GCNSchedStrategy::initCandidate(SchedCandidate &Cand, SUnit *SU,
}
// Register pressure is considered 'CRITICAL' if it is approaching a value
- // that would reduce the wave occupancy for the execution unit. When
+ // that would reduce the wave occupancy for the execution unit. When
// register pressure is 'CRITICAL', increasing SGPR and VGPR pressure both
// has the same cost, so we don't need to prefer one over the other.
@@ -848,8 +848,8 @@ GCNMaxMemoryClauseSchedStrategy::GCNMaxMemoryClauseSchedStrategy(
/// GCNMaxMemoryClauseSchedStrategy tries best to clause memory instructions as
/// much as possible. This is achieved by:
-// 1. Prioritize clustered operations before stall latency heuristic.
-// 2. Prioritize long-latency-load before stall latency heuristic.
+/// 1. Prioritize clustered operations before stall latency heuristic.
+/// 2. Prioritize long-latency-load before stall latency heuristic.
///
/// \param Cand provides the policy and current best candidate.
/// \param TryCand refers to the next SUnit candidate, otherwise uninitialized.
@@ -1445,7 +1445,7 @@ Printable PreRARematStage::ScoredRemat::print() const {
bool PreRARematStage::initGCNSchedStage() {
// FIXME: This pass will invalidate cached BBLiveInMap and MBBLiveIns for
- // regions inbetween the defs and region we sinked the def to. Will need to be
+ // regions in between the defs and region we sunk the def to. Will need to be
// fixed if there is another pass after this pass.
assert(!S.hasNextStage());
@@ -1454,7 +1454,7 @@ bool PreRARematStage::initGCNSchedStage() {
// Maps all MIs (except lone terminators, which are not part of any region) to
// their parent region. Non-lone terminators are considered part of the region
- // they delimitate.
+ // they delimit.
DenseMap<MachineInstr *, unsigned> MIRegion(MF.getInstructionCount());
// Before performing any IR modification record the parent region of each MI
@@ -1560,7 +1560,7 @@ bool PreRARematStage::initGCNSchedStage() {
});
// Rematerialize registers in successive rounds until all RP targets are
- // satisifed or until we run out of rematerialization candidates.
+ // satisfied or until we run out of rematerialization candidates.
BitVector RecomputeRP(DAG.Regions.size());
for (;;) {
RecomputeRP.reset();
@@ -1820,9 +1820,9 @@ bool UnclusteredHighRPStage::initGCNRegion() {
bool ClusteredLowOccStage::initGCNRegion() {
// We may need to reschedule this region if it wasn't rescheduled in the last
// stage, or if we found it was testing critical register pressure limits in
- // the unclustered reschedule stage. The later is because we may not have been
- // able to raise the min occupancy in the previous stage so the region may be
- // overly constrained even if it was already rescheduled.
+ // the unclustered reschedule stage. The latter is because we may not have
+ // been able to raise the min occupancy in the previous stage so the region
+ // may be overly constrained even if it was already rescheduled.
if (!DAG.RegionsWithHighRP[RegionIdx])
return false;
@@ -2266,10 +2266,8 @@ void RewriteMFMAFormStage::resetRewriteCandsToVGPR(
// Have to get src types separately since subregs may cause C and D
// registers to be different types even though the actual operand is
// the same size.
- const TargetRegisterClass *AUseRC =
- DAG.MRI.getRegClass(Src2->getReg());
- const TargetRegisterClass *VUseRC =
- SRI->getEquivalentVGPRClass(AUseRC);
+ const TargetRegisterClass *AUseRC = DAG.MRI.getRegClass(Src2->getReg());
+ const TargetRegisterClass *VUseRC = SRI->getEquivalentVGPRClass(AUseRC);
DAG.MRI.setRegClass(Src2->getReg(), VUseRC);
}
}
@@ -2387,7 +2385,7 @@ int64_t RewriteMFMAFormStage::getRewriteCost(
// Reset the classes that were changed to AGPR for better RB analysis.
// We must do rewriting after copy-insertion, as some defs of the register
- // may require VGPR. Additionally, if we bail out and don't perform the
+ // may require VGPR. Additionally, if we bail out and don't perform the
// rewrite then these need to be restored anyway.
for (unsigned Region = 0; Region < DAG.Regions.size(); Region++) {
if (!RegionsWithExcessArchVGPR[Region])
@@ -2660,7 +2658,7 @@ bool RewriteMFMAFormStage::rewrite(
// If none exists, create a copy from this reaching def.
// We may have inserted a copy already in an earlier iteration.
for (MachineInstr *RD : DstUseDefsReplace) {
- // Do not create reundant copies.
+ // Do not create redundant copies.
if (ReachingDefCopyMap[DstReg].insert(RD).second) {
MachineInstrBuilder VGPRCopy =
BuildMI(*RD->getParent(), std::next(RD->getIterator()),
@@ -2870,7 +2868,7 @@ bool PreRARematStage::collectRematRegs(
// regions containing rematerializable instructions.
DAG.RegionLiveOuts.buildLiveRegMap();
- // Set of registers already marked for potential remterialization; used to
+ // Set of registers already marked for potential rematerialization; used to
// avoid rematerialization chains.
SmallSet<Register, 4> MarkedRegs;
auto IsMarkedForRemat = [&MarkedRegs](const MachineOperand &MO) -> bool {
@@ -2915,7 +2913,7 @@ bool PreRARematStage::collectRematRegs(
llvm::any_of(DefMI.operands(), IsMarkedForRemat))
continue;
- // Do not rematerialize an instruction it it uses registers that aren't
+ // Do not rematerialize an instruction if it uses registers that aren't
// available at its use. This ensures that we are not extending any live
// range while rematerializing.
SlotIndex UseIdx = DAG.LIS->getInstructionIndex(*UseMI).getRegSlot(true);
@@ -3109,7 +3107,7 @@ PreRARematStage::ScoredRemat::rematerialize(GCNScheduleDAGMILive &DAG) const {
}
void PreRARematStage::commitRematerializations() const {
- REMAT_DEBUG(dbgs() << "Commiting all rematerializations\n");
+ REMAT_DEBUG(dbgs() << "Committing all rematerializations\n");
for (const RollbackInfo &Rollback : Rollbacks)
DAG.deleteMI(Rollback.Remat->DefRegion, Rollback.Remat->DefMI);
}
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
index 684267d381478..759a88ccfbc8b 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
@@ -41,7 +41,7 @@ enum class GCNSchedStageID : unsigned {
raw_ostream &operator<<(raw_ostream &OS, const GCNSchedStageID &StageID);
#endif
-/// This is a minimal scheduler strategy. The main difference between this
+/// This is a minimal scheduler strategy. The main difference between this
/// and the GenericScheduler is that GCNSchedStrategy uses different
/// heuristics to determine excess/critical pressure sets.
class GCNSchedStrategy : public GenericScheduler {
@@ -104,7 +104,7 @@ class GCNSchedStrategy : public GenericScheduler {
// GCN RP Tracker for top-down scheduling
mutable GCNDownwardRPTracker DownwardTracker;
- // GCN RP Tracker for botttom-up scheduling
+ // GCN RP Tracker for bottom-up scheduling
mutable GCNUpwardRPTracker UpwardTracker;
bool UseGCNTrackers = false;
@@ -312,7 +312,7 @@ class GCNScheduleDAGMILive final : public ScheduleDAGMILive {
// The live out registers per region. These are internally stored as a map of
// the initial last region instruction to region live out registers, but can
- // be retreived with the regionIdx by calls to getLiveRegsForRegionIdx.
+ // be retrieved with the regionIdx by calls to getLiveRegsForRegionIdx.
RegionPressureMap RegionLiveOuts;
// Return current region pressure.
@@ -394,7 +394,7 @@ class GCNSchedStage {
// Check result of scheduling.
void checkScheduling();
- // computes the given schedule virtual execution time in clocks
+ // Computes the given schedule virtual execution time in clocks
ScheduleMetrics getScheduleMetrics(const std::vector<SUnit> &InputSchedule);
ScheduleMetrics getScheduleMetrics(const GCNScheduleDAGMILive &DAG);
unsigned computeSUnitReadyCycle(const SUnit &SU, unsigned CurrCycle,
@@ -447,7 +447,7 @@ class RewriteMFMAFormStage : public GCNSchedStage {
/// Do a speculative rewrite and collect copy locations. The speculative
/// rewrite allows us to calculate the RP of the code after the rewrite, and
/// the copy locations allow us to calculate the total cost of copies required
- /// for the rewrite. Stores the rewritten instructions in \p RewriteCands ,
+ /// for the rewrite. Stores the rewritten instructions in \p RewriteCands,
/// the copy locations for uses (of the MFMA result) in \p CopyForUse and the
/// copy locations for defs (of the MFMA operands) in \p CopyForDef
bool
@@ -464,10 +464,10 @@ class RewriteMFMAFormStage : public GCNSchedStage {
const SmallPtrSetImpl<MachineInstr *> &CopyForDef);
/// Do the final rewrite on \p RewriteCands and insert any needed copies.
- bool
- rewrite(ArrayRef<std::pair<MachineInstr *, unsigned>> RewriteCands);
- /// Resets all rewrite candidates in \p Cands back to their original VGPR
- /// opcodes and register classes.
+ bool rewrite(ArrayRef<std::pair<MachineInstr *, unsigned>> RewriteCands);
+
+ /// Resets all rewrite candidates in \p RewriteCands back to their original
+ /// VGPR opcodes and register classes.
void resetRewriteCandsToVGPR(
ArrayRef<std::pair<MachineInstr *, unsigned>> RewriteCands);
@@ -541,7 +541,7 @@ class ClusteredLowOccStage : public GCNSchedStage {
/// 2. The single defining instruction is either deemed rematerializable by the
/// target-independent logic, or if not, has no non-constant and
/// non-ignorable physical register use.
-/// 3 The register has no virtual register use whose live range would be
+/// 3. The register has no virtual register use whose live range would be
/// extended by the rematerialization.
/// 4. The register has a single non-debug user in a different region from its
/// defining region.
@@ -570,8 +570,8 @@ class PreRARematStage : public GCNSchedStage {
};
/// A scored rematerialization candidate. Higher scores indicate more
- /// beneficial rematerializations. A null score indicate the rematerialization
- /// is not helpful to reduce RP in target regions.
+ /// beneficial rematerializations. A null score indicates the
+ /// rematerialization is not helpful to reduce RP in target regions.
struct ScoredRemat {
/// The rematerializable register under consideration.
RematReg *Remat;
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-check-half-rewrite.mir b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-check-half-rewrite.mir
index 3a73e82395b9b..3ff5e7d5a7694 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-check-half-rewrite.mir
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-check-half-rewrite.mir
@@ -66,11 +66,10 @@ body: |
%6:vreg_1024 = IMPLICIT_DEF
%7:vreg_1024 = IMPLICIT_DEF
%8:vreg_1024 = IMPLICIT_DEF
-
- %9:av_128_align2 = IMPLICIT_DEF
- %10:av_128_align2 = IMPLICIT_DEF
+ %9:av_128_align2 = IMPLICIT_DEF
+ %10:av_128_align2 = IMPLICIT_DEF
%11:vreg_64_align2 = IMPLICIT_DEF
- %12:vgpr_32 = IMPLICIT_DEF
+ %12:vgpr_32 = IMPLICIT_DEF
SCHED_BARRIER 0
More information about the llvm-commits
mailing list