[llvm] [AMDGPU] Fix RewriteMFMAFormSchedStage (PR #193429)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Apr 22 20:52:06 PDT 2026
https://github.com/xgxanq updated https://github.com/llvm/llvm-project/pull/193429
>From 9b182e0ffa19bbf53d942ac8767217cb8aead197 Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Wed, 22 Apr 2026 07:21:32 +0000
Subject: [PATCH 1/2] [AMDGPU] Fix RewriteMFMAFormSchedStage
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 114 ++++---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.h | 5 +
.../rewrite-mfma-form-ReachingUseTracker.ll | 213 ++++++++++++++
.../AMDGPU/rewrite-mfma-form-case2-case3.ll | 147 +++++++++
.../rewrite-mfma-form-check-half-rewrite.mir | 110 +++++++
.../rewrite-mfma-form-v7slice-pattern.ll | 278 ++++++++++++++++++
6 files changed, 822 insertions(+), 45 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-ReachingUseTracker.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-case2-case3.ll
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..980beb38cd46d 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -2250,6 +2250,31 @@ void GCNSchedStage::modifyRegionSchedule(unsigned RegionIdx,
DAG.Regions[RegionIdx].first = MIOrder.front();
}
+void RewriteMFMAFormStage::resetRewriteCandsToVGPR(
+ const std::vector<std::pair<MachineInstr *, unsigned>> &RewriteCands) {
+ 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);
+ 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 +2297,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 +2322,6 @@ bool RewriteMFMAFormStage::initHeuristics(
}
}
- MachineOperand &Dst = MI.getOperand(0);
SmallVector<MachineOperand *, 8> DstReachingUses;
findReachingUses(&MI, DAG.LIS, DstReachingUses);
@@ -2354,6 +2386,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 +2427,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,30 +2467,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
- // 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;
}
@@ -2534,9 +2549,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 +2614,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 +2684,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 +2701,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 +2817,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..258010a071704 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
@@ -467,6 +467,11 @@ class RewriteMFMAFormStage : public GCNSchedStage {
bool
rewrite(const std::vector<std::pair<MachineInstr *, unsigned>> &RewriteCands);
+ /// Resets all rewrite candidates in \p Cands back to their original VGPR
+ /// opcodes and register classes.
+ void resetRewriteCandsToVGPR(
+ const std::vector<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-ReachingUseTracker.ll b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-ReachingUseTracker.ll
new file mode 100644
index 0000000000000..13bc16aad14a4
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-ReachingUseTracker.ll
@@ -0,0 +1,213 @@
+;
+; Verify the machine-scheduler-level rewrite() changes:
+; 1. All 9 MFMAs: opcode _vgprcd_e64 -> _e64, class vreg -> areg (AGPR form).
+; 2. Case 3 COPY: inserted after zeroinitializer def in loop.body (entry edge),
+; converts VGPR acc -> AGPR acc before MFMA (src2/dst reclassification).
+; 3. Case 2 COPY: inserted at exit block entry (ReachingUseTracker),
+; converts AGPR MFMA result -> VGPR before fptrunc (V_CVT_F16_F32).
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a \
+; RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+; RUN: < %s | FileCheck %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a \
+; RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+; RUN: -stop-after=machine-scheduler \
+; RUN: < %s | FileCheck %s --check-prefix=MIR
+;
+; CHECK-LABEL: test_ReachingUseTracker_crossblock_use:
+; CHECK-COUNT-9: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], {{.*}}, {{.*}}, a[{{[0-9:]+}}]
+;
+; MIR-LABEL: name: test_ReachingUseTracker_crossblock_use
+; MIR-LABEL: bb.0.entry:
+;
+; Case 3: 9 x (areg_128_align2 = COPY vreg) in entry block —
+; rewrite() inserts these after the zeroinitializer reaching-defs of MFMA src2.
+; Converts VGPR accumulator -> AGPR before first loop iteration.
+; MIR-COUNT-9: areg_128_align2 = COPY
+;
+; MIR-LABEL: bb.1.loop.body:
+;
+; All 9 MFMAs rewritten: _vgprcd_e64 (VGPR C/D) -> _e64 (AGPR), dst is areg.
+; MIR-COUNT-9: areg_128_align2 = V_MFMA_F32_4X4X2BF16_e64
+; MIR-NOT: V_MFMA_F32_4X4X2BF16_vgprcd_e64
+;
+; MIR-LABEL: bb.2.exit:
+;
+; Case 2: 9 x (vreg_128_align2 = COPY areg) at exit block start —
+; rewrite() inserts these via ReachingUseTracker.
+; MFMA dst (AGPR) -> VGPR before fptrunc (V_CVT_F16_F32).
+; MIR-COUNT-9: vreg_128_align2 = COPY
+;
+; fptrunc lowering: V_CVT_F16_F32 operates on the new VGPR copies.
+; MIR: V_CVT_F16_F32_e32
+
+; Test: triggers RewriteMFMAFormStage::rewrite() with getRewriteCost() < 0
+; and exercises the cross-block ReachingUseTracker path.
+;
+; Pressure design (gfx90a, separate AGPR/VGPR files):
+; 8 loop-carried <32 x float> phi nodes (c0..c7) = 8*32=256 VGPRs non-MFMA.
+; Carrier back-edge: self-insertelement (no dep on MFMA results).
+; Carrier use in loop.body: store one element -> keeps carriers live.
+; GPU scheduler issues MFMAs first (long latency), then fills with stores.
+; Carriers naturally live across MFMAs: peak = 256+36 = 292 > 256 VGPRs.
+; After rewrite: ArchVGPR=256 (carriers), AGPR=36 < 256 -> no AGPR spill.
+; SpillBenefit = -(292-256)*2 * LoopFreq = -72*LoopFreq << 0
+; MFMA results used ONLY in exit block -> UseFreq~0 -> CopyCost~36.
+; getRewriteCost() << 0 -> rewrite() called.
+;
+; Cross-block use (ReachingUseTracker):
+; MFMAs in loop.body; fptrunc+store in exit block.
+; RUBlock(exit) != MI->getParent()(loop.body) -> ReachingUseTracker.
+
+declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16>, <2 x i16>, <4 x float>, i32 immarg, i32 immarg, i32 immarg)
+
+define amdgpu_kernel void @test_ReachingUseTracker_crossblock_use(
+ ptr addrspace(1) %out,
+ <2 x i16> %a,
+ <2 x i16> %b,
+ i32 %n) #0 {
+entry:
+ br label %loop.body
+
+loop.body:
+ %i = phi i32 [ 0, %entry ], [ %i.next, %loop.body ]
+ ; 8 loop-carried <32 x float>: 8*32=256 VGPRs non-MFMA pressure.
+ ; Back-edge: self-insert (no dep on MFMA results, never DCE'd due to store use).
+ %c0 = phi <32 x float> [ zeroinitializer, %entry ], [ %c0n, %loop.body ]
+ %c1 = phi <32 x float> [ zeroinitializer, %entry ], [ %c1n, %loop.body ]
+ %c2 = phi <32 x float> [ zeroinitializer, %entry ], [ %c2n, %loop.body ]
+ %c3 = phi <32 x float> [ zeroinitializer, %entry ], [ %c3n, %loop.body ]
+ %c4 = phi <32 x float> [ zeroinitializer, %entry ], [ %c4n, %loop.body ]
+ %c5 = phi <32 x float> [ zeroinitializer, %entry ], [ %c5n, %loop.body ]
+ %c6 = phi <32 x float> [ zeroinitializer, %entry ], [ %c6n, %loop.body ]
+ %c7 = phi <32 x float> [ zeroinitializer, %entry ], [ %c7n, %loop.body ]
+ ; 9 MFMA accumulators: 9*4=36 VGPRs.
+ %acc0 = phi <4 x float> [ zeroinitializer, %entry ], [ %r0, %loop.body ]
+ %acc1 = phi <4 x float> [ zeroinitializer, %entry ], [ %r1, %loop.body ]
+ %acc2 = phi <4 x float> [ zeroinitializer, %entry ], [ %r2, %loop.body ]
+ %acc3 = phi <4 x float> [ zeroinitializer, %entry ], [ %r3, %loop.body ]
+ %acc4 = phi <4 x float> [ zeroinitializer, %entry ], [ %r4, %loop.body ]
+ %acc5 = phi <4 x float> [ zeroinitializer, %entry ], [ %r5, %loop.body ]
+ %acc6 = phi <4 x float> [ zeroinitializer, %entry ], [ %r6, %loop.body ]
+ %acc7 = phi <4 x float> [ zeroinitializer, %entry ], [ %r7, %loop.body ]
+ %acc8 = phi <4 x float> [ zeroinitializer, %entry ], [ %r8, %loop.body ]
+
+ ; 9 MFMAs. Results only used in exit (CopyForUse exit only, UseFreq~0).
+ %r0 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc0, i32 0, i32 0, i32 0)
+ %r1 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc1, i32 0, i32 0, i32 0)
+ %r2 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc2, i32 0, i32 0, i32 0)
+ %r3 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc3, i32 0, i32 0, i32 0)
+ %r4 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc4, i32 0, i32 0, i32 0)
+ %r5 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc5, i32 0, i32 0, i32 0)
+ %r6 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc6, i32 0, i32 0, i32 0)
+ %r7 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc7, i32 0, i32 0, i32 0)
+ %r8 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc8, i32 0, i32 0, i32 0)
+
+ ; Variable index prevents folding of carrier self-inserts.
+ %eidx = and i32 %i, 31
+ ; Carrier self-insert: no dep on MFMA results. GPU scheduler fills MFMA
+ ; latency with these, naturally scheduling them AFTER MFMAs -> carriers live
+ ; across all MFMAs -> peak pressure 256+36=292 VGPRs.
+ %c0e = extractelement <32 x float> %c0, i32 0
+ %c1e = extractelement <32 x float> %c1, i32 0
+ %c2e = extractelement <32 x float> %c2, i32 0
+ %c3e = extractelement <32 x float> %c3, i32 0
+ %c4e = extractelement <32 x float> %c4, i32 0
+ %c5e = extractelement <32 x float> %c5, i32 0
+ %c6e = extractelement <32 x float> %c6, i32 0
+ %c7e = extractelement <32 x float> %c7, i32 0
+ %c0n = insertelement <32 x float> %c0, float %c0e, i32 %eidx
+ %c1n = insertelement <32 x float> %c1, float %c1e, i32 %eidx
+ %c2n = insertelement <32 x float> %c2, float %c2e, i32 %eidx
+ %c3n = insertelement <32 x float> %c3, float %c3e, i32 %eidx
+ %c4n = insertelement <32 x float> %c4, float %c4e, i32 %eidx
+ %c5n = insertelement <32 x float> %c5, float %c5e, i32 %eidx
+ %c6n = insertelement <32 x float> %c6, float %c6e, i32 %eidx
+ %c7n = insertelement <32 x float> %c7, float %c7e, i32 %eidx
+
+ ; Store one carrier element (keeps carriers live, prevents DCE).
+ %csum = fadd float %c0e, %c1e
+ store float %csum, ptr addrspace(1) %out, align 4
+
+ %i.next = add i32 %i, 1
+ %cond = icmp eq i32 %i.next, %n
+ br i1 %cond, label %exit, label %loop.body
+
+exit:
+ ; Cross-block non-MAI uses: MFMAs in loop.body, fptrunc uses in exit.
+ ; -> ReachingUseTracker.
+ %e0_0 = extractelement <4 x float> %r0, i32 0
+ %e1_0 = extractelement <4 x float> %r0, i32 1
+ %v0a = insertelement <2 x float> poison, float %e0_0, i32 0
+ %v0 = insertelement <2 x float> %v0a, float %e1_0, i32 1
+ %h0 = fptrunc <2 x float> %v0 to <2 x half>
+ %p_h0 = getelementptr i16, ptr addrspace(1) %out, i32 2
+ store <2 x half> %h0, ptr addrspace(1) %p_h0, align 2
+
+ %e0_1 = extractelement <4 x float> %r1, i32 0
+ %e1_1 = extractelement <4 x float> %r1, i32 1
+ %v1a = insertelement <2 x float> poison, float %e0_1, i32 0
+ %v1 = insertelement <2 x float> %v1a, float %e1_1, i32 1
+ %h1 = fptrunc <2 x float> %v1 to <2 x half>
+ %p_h1 = getelementptr i16, ptr addrspace(1) %out, i32 4
+ store <2 x half> %h1, ptr addrspace(1) %p_h1, align 2
+
+ %e0_2 = extractelement <4 x float> %r2, i32 0
+ %e1_2 = extractelement <4 x float> %r2, i32 1
+ %v2a = insertelement <2 x float> poison, float %e0_2, i32 0
+ %v2 = insertelement <2 x float> %v2a, float %e1_2, i32 1
+ %h2 = fptrunc <2 x float> %v2 to <2 x half>
+ %p_h2 = getelementptr i16, ptr addrspace(1) %out, i32 6
+ store <2 x half> %h2, ptr addrspace(1) %p_h2, align 2
+
+ %e0_3 = extractelement <4 x float> %r3, i32 0
+ %e1_3 = extractelement <4 x float> %r3, i32 1
+ %v3a = insertelement <2 x float> poison, float %e0_3, i32 0
+ %v3 = insertelement <2 x float> %v3a, float %e1_3, i32 1
+ %h3 = fptrunc <2 x float> %v3 to <2 x half>
+ %p_h3 = getelementptr i16, ptr addrspace(1) %out, i32 8
+ store <2 x half> %h3, ptr addrspace(1) %p_h3, align 2
+
+ %e0_4 = extractelement <4 x float> %r4, i32 0
+ %e1_4 = extractelement <4 x float> %r4, i32 1
+ %v4a = insertelement <2 x float> poison, float %e0_4, i32 0
+ %v4 = insertelement <2 x float> %v4a, float %e1_4, i32 1
+ %h4 = fptrunc <2 x float> %v4 to <2 x half>
+ %p_h4 = getelementptr i16, ptr addrspace(1) %out, i32 10
+ store <2 x half> %h4, ptr addrspace(1) %p_h4, align 2
+
+ %e0_5 = extractelement <4 x float> %r5, i32 0
+ %e1_5 = extractelement <4 x float> %r5, i32 1
+ %v5a = insertelement <2 x float> poison, float %e0_5, i32 0
+ %v5 = insertelement <2 x float> %v5a, float %e1_5, i32 1
+ %h5 = fptrunc <2 x float> %v5 to <2 x half>
+ %p_h5 = getelementptr i16, ptr addrspace(1) %out, i32 12
+ store <2 x half> %h5, ptr addrspace(1) %p_h5, align 2
+
+ %e0_6 = extractelement <4 x float> %r6, i32 0
+ %e1_6 = extractelement <4 x float> %r6, i32 1
+ %v6a = insertelement <2 x float> poison, float %e0_6, i32 0
+ %v6 = insertelement <2 x float> %v6a, float %e1_6, i32 1
+ %h6 = fptrunc <2 x float> %v6 to <2 x half>
+ %p_h6 = getelementptr i16, ptr addrspace(1) %out, i32 14
+ store <2 x half> %h6, ptr addrspace(1) %p_h6, align 2
+
+ %e0_7 = extractelement <4 x float> %r7, i32 0
+ %e1_7 = extractelement <4 x float> %r7, i32 1
+ %v7a = insertelement <2 x float> poison, float %e0_7, i32 0
+ %v7 = insertelement <2 x float> %v7a, float %e1_7, i32 1
+ %h7 = fptrunc <2 x float> %v7 to <2 x half>
+ %p_h7 = getelementptr i16, ptr addrspace(1) %out, i32 16
+ store <2 x half> %h7, ptr addrspace(1) %p_h7, align 2
+
+ %e0_8 = extractelement <4 x float> %r8, i32 0
+ %e1_8 = extractelement <4 x float> %r8, i32 1
+ %v8a = insertelement <2 x float> poison, float %e0_8, i32 0
+ %v8 = insertelement <2 x float> %v8a, float %e1_8, i32 1
+ %h8 = fptrunc <2 x float> %v8 to <2 x half>
+ %p_h8 = getelementptr i16, ptr addrspace(1) %out, i32 18
+ store <2 x half> %h8, ptr addrspace(1) %p_h8, align 2
+
+ ret void
+}
+
+attributes #0 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" }
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-case2-case3.ll b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-case2-case3.ll
new file mode 100644
index 0000000000000..a698764417004
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-case2-case3.ll
@@ -0,0 +1,147 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a \
+; RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+; RUN: < %s | FileCheck %s
+;
+; Verify MIR state BEFORE machine-scheduler (rewrite not yet applied):
+; All 9 MFMAs use VGPR form: vreg_128_align2 / V_MFMA_F32_4X4X2BF16_vgprcd_e64.
+; No AGPR COPYs inserted yet.
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a \
+; RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+; RUN: -stop-before=machine-scheduler \
+; RUN: < %s | FileCheck %s --check-prefix=BEFORE
+;
+; Verify MIR state AFTER machine-scheduler (rewrite applied):
+; Case 3: 9 x (areg_128_align2 = COPY vreg) in bb.0.entry.
+; loop.body: all 9 MFMAs rewritten to V_MFMA_F32_4X4X2BF16_e64 with areg dst.
+; Case 2: vreg_128_align2 = COPY areg in bb.2.exit.
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a \
+; RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+; RUN: -stop-after=machine-scheduler \
+; RUN: < %s | FileCheck %s --check-prefix=MIR
+;
+; BEFORE-LABEL: name: test_case2_case3
+; BEFORE-LABEL: bb.1.loop.body:
+; BEFORE-COUNT-9: vreg_128_align2 = V_MFMA_F32_4X4X2BF16_vgprcd_e64
+; BEFORE-NOT: V_MFMA_F32_4X4X2BF16_e64
+; BEFORE-LABEL: bb.2.exit:
+; BEFORE-NOT: areg_128_align2 = COPY
+;
+; MIR-LABEL: name: test_case2_case3
+; MIR-LABEL: bb.0.entry:
+;
+; Case 3: 9 x (areg_128_align2 = COPY vreg) inserted in entry.
+; Converts VGPR zeroinit accumulator -> AGPR before first loop iteration.
+; MIR-COUNT-9: areg_128_align2 = COPY
+;
+; MIR-LABEL: bb.1.loop.body:
+;
+; All 9 MFMAs rewritten: _vgprcd_e64 (VGPR C/D) -> _e64 (AGPR), dst is areg.
+; MIR-COUNT-9: areg_128_align2 = V_MFMA_F32_4X4X2BF16_e64
+; MIR-NOT: V_MFMA_F32_4X4X2BF16_vgprcd_e64
+;
+; MIR-LABEL: bb.2.exit:
+;
+; Case 2: vreg_128_align2 = COPY areg inserted at exit block start.
+; Converts AGPR MFMA result -> VGPR before extractelement use.
+; MIR: vreg_128_align2 = COPY
+
+; Test: RewriteMFMAFormStage — Case 2 + Case 3
+;
+; Case 2: MFMA dst has a non-MAI USE in a different block (exit).
+; Case 3: At that use point, %acc has a non-MAI reaching def (zeroinitializer
+; from entry, lowered to v_mov_b32).
+;
+; Pressure design (gfx90a, minWavesPerEU=1):
+; 8 loop-carried <32 x float> phi nodes (c0..c7) = 8*32 = 256 VGPRs non-MFMA.
+; Variable-index insertelement prevents carrier DCE and keeps all carriers
+; live across the MFMA region.
+; 9 MFMA accumulators (mfma.f32.4x4x2bf16, dst=<4 x float>) = 9*4 = 36 VGPRs.
+; Peak ArchVGPR = 256+36 = 292 > 256 (gfx90a limit) -> rewrite fires.
+;
+; Expected output (rewrite fires, cost <= 0):
+; Case 3: v_accvgpr_write (VGPR->AGPR) in entry for each MFMA zeroinit.
+; Case 2: v_accvgpr_read (AGPR->VGPR) in exit before extractelement uses.
+
+declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16>, <2 x i16>, <4 x float>, i32 immarg, i32 immarg, i32 immarg)
+
+define amdgpu_kernel void @test_case2_case3(
+; CHECK-COUNT-9: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], {{.*}}, {{.*}}, a[{{[0-9:]+}}]
+ ptr addrspace(1) %out,
+ <2 x i16> %a,
+ <2 x i16> %b,
+ i32 %n) #0 {
+entry:
+ br label %loop.body
+
+loop.body:
+ %i = phi i32 [ 0, %entry ], [ %i.next, %loop.body ]
+ ; 8 loop-carried <32 x float>: 8*32=256 VGPRs non-MFMA pressure.
+ %c0 = phi <32 x float> [ zeroinitializer, %entry ], [ %c0n, %loop.body ]
+ %c1 = phi <32 x float> [ zeroinitializer, %entry ], [ %c1n, %loop.body ]
+ %c2 = phi <32 x float> [ zeroinitializer, %entry ], [ %c2n, %loop.body ]
+ %c3 = phi <32 x float> [ zeroinitializer, %entry ], [ %c3n, %loop.body ]
+ %c4 = phi <32 x float> [ zeroinitializer, %entry ], [ %c4n, %loop.body ]
+ %c5 = phi <32 x float> [ zeroinitializer, %entry ], [ %c5n, %loop.body ]
+ %c6 = phi <32 x float> [ zeroinitializer, %entry ], [ %c6n, %loop.body ]
+ %c7 = phi <32 x float> [ zeroinitializer, %entry ], [ %c7n, %loop.body ]
+ ; 9 MFMA accumulators: 9*4=36 VGPRs. Case 3: entry zeroinitializer -> AGPR.
+ %acc0 = phi <4 x float> [ zeroinitializer, %entry ], [ %r0, %loop.body ]
+ %acc1 = phi <4 x float> [ zeroinitializer, %entry ], [ %r1, %loop.body ]
+ %acc2 = phi <4 x float> [ zeroinitializer, %entry ], [ %r2, %loop.body ]
+ %acc3 = phi <4 x float> [ zeroinitializer, %entry ], [ %r3, %loop.body ]
+ %acc4 = phi <4 x float> [ zeroinitializer, %entry ], [ %r4, %loop.body ]
+ %acc5 = phi <4 x float> [ zeroinitializer, %entry ], [ %r5, %loop.body ]
+ %acc6 = phi <4 x float> [ zeroinitializer, %entry ], [ %r6, %loop.body ]
+ %acc7 = phi <4 x float> [ zeroinitializer, %entry ], [ %r7, %loop.body ]
+ %acc8 = phi <4 x float> [ zeroinitializer, %entry ], [ %r8, %loop.body ]
+
+ ; 9 MFMAs — results only used in exit (Case 2: cross-block non-MAI use).
+ %r0 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc0, i32 0, i32 0, i32 0)
+ %r1 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc1, i32 0, i32 0, i32 0)
+ %r2 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc2, i32 0, i32 0, i32 0)
+ %r3 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc3, i32 0, i32 0, i32 0)
+ %r4 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc4, i32 0, i32 0, i32 0)
+ %r5 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc5, i32 0, i32 0, i32 0)
+ %r6 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc6, i32 0, i32 0, i32 0)
+ %r7 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc7, i32 0, i32 0, i32 0)
+ %r8 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc8, i32 0, i32 0, i32 0)
+
+ ; Variable index prevents folding: keeps carriers live across all MFMAs.
+ %eidx = and i32 %i, 31
+ %c0e = extractelement <32 x float> %c0, i32 0
+ %c1e = extractelement <32 x float> %c1, i32 0
+ %c2e = extractelement <32 x float> %c2, i32 0
+ %c3e = extractelement <32 x float> %c3, i32 0
+ %c4e = extractelement <32 x float> %c4, i32 0
+ %c5e = extractelement <32 x float> %c5, i32 0
+ %c6e = extractelement <32 x float> %c6, i32 0
+ %c7e = extractelement <32 x float> %c7, i32 0
+ %c0n = insertelement <32 x float> %c0, float %c0e, i32 %eidx
+ %c1n = insertelement <32 x float> %c1, float %c1e, i32 %eidx
+ %c2n = insertelement <32 x float> %c2, float %c2e, i32 %eidx
+ %c3n = insertelement <32 x float> %c3, float %c3e, i32 %eidx
+ %c4n = insertelement <32 x float> %c4, float %c4e, i32 %eidx
+ %c5n = insertelement <32 x float> %c5, float %c5e, i32 %eidx
+ %c6n = insertelement <32 x float> %c6, float %c6e, i32 %eidx
+ %c7n = insertelement <32 x float> %c7, float %c7e, i32 %eidx
+
+ %csum = fadd float %c0e, %c1e
+ store float %csum, ptr addrspace(1) %out, align 4
+
+ %i.next = add i32 %i, 1
+ %cond = icmp eq i32 %i.next, %n
+ br i1 %cond, label %exit, label %loop.body
+
+exit:
+ ; Case 2: non-MAI uses of MFMA dst in exit block (different block from loop.body).
+ ; Case 3: reaching def of each %acc in entry is zeroinitializer (non-MAI).
+ ; Expected: v_accvgpr_read (AGPR->VGPR) inserted before these extractelements.
+ %e0 = extractelement <4 x float> %r0, i32 0
+ %e1 = extractelement <4 x float> %r1, i32 0
+ store float %e0, ptr addrspace(1) %out, align 4
+ %p1 = getelementptr float, ptr addrspace(1) %out, i32 1
+ store float %e1, ptr addrspace(1) %p1, align 4
+ ret void
+}
+
+attributes #0 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" }
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..0f78ee7ee163d
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-check-half-rewrite.mir
@@ -0,0 +1,110 @@
+# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 5
+# 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: Bug 2 in RewriteMFMAFormStage — physical register in MFMA src2.
+#
+# Root cause (GCNSchedStrategy.cpp):
+# rewrite() processes RewriteCands in iteration order. For each MFMA:
+# MI->setDesc(ReplacementOp); // opcode mutated
+# if (!Src2Reg.isVirtual()) return false; // Bug 2: mid-loop bail-out
+# A return false here leaves all previously-processed MFMAs in AGPR form
+# while subsequent MFMAs are partially mutated — inconsistent MIR state.
+#
+# Trigger:
+# MFMA_B uses $vgpr4_vgpr5_vgpr6_vgpr7 (physical VGPR) as src2.
+# VISrc_128_f32 maps to physical class VGPR_128, so MIR verification passes.
+# Without the fix: initHeuristics calls findReachingDefs($vgpr4) ->
+# LIS->getInterval($vgpr4) -> virtRegIndex() asserts "Not a virtual register".
+#
+# Fix (initHeuristics pre-validation):
+# Before calling findReachingDefs/Uses or getRegClass/setRegClass, check that
+# both dst and src2 (if reg) are virtual. If not, skip this MFMA candidate.
+# With the fix: MFMA_B is skipped -> not added to RewriteCands -> rewrite()
+# never sees a non-virtual register -> no crash, no half-rewritten state.
+# Since the rewrite cost changes, MFMA_A also remains in VGPR form.
+#
+# Pressure setup: 9 x vreg_1024 = 288 VGPRs > 256 (gfx950 ArchVGPR limit),
+# satisfying RegionsWithExcessArchVGPR in initGCNSchedStage().
+
+--- |
+ 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: |
+ ; CHECK-LABEL: name: test_bug2_physreg_src2
+ ; CHECK: bb.0:
+ ; CHECK: liveins: $vgpr0, $sgpr4_sgpr5, $vgpr4, $vgpr5, $vgpr6, $vgpr7
+ ;
+ ; After the fix, MFMA_B (with physical src2) is skipped in initHeuristics.
+ ; The combined rewrite is not beneficial, so MFMA_A also remains in VGPR form.
+ ; Both MFMAs must stay as _vgprcd_e64 (not _agprcd_e64).
+ ;
+ ; CHECK: V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64
+ ; CHECK: V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64
+ ; CHECK-NOT: V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_agprcd_e64
+ 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
+
+ ; MFMA src0/src1 inputs.
+ %10:av_128_align2 = IMPLICIT_DEF
+ %11:av_128_align2 = IMPLICIT_DEF
+ %12:vreg_64_align2 = IMPLICIT_DEF
+ %13:vgpr_32 = IMPLICIT_DEF
+
+ SCHED_BARRIER 0
+
+ ; MFMA_A: all virtual registers — normal rewrite candidate.
+ ; Without the fix, this would be fully rewritten to AGPR form before
+ ; MFMA_B is processed, causing the half-rewritten inconsistent state.
+ %acc_a:vreg_128_align2 = IMPLICIT_DEF
+ %res_a:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 %10:av_128_align2, %11:av_128_align2, %acc_a:vreg_128_align2, 4, 4, %12.sub0:vreg_64_align2, %13:vgpr_32, 0, 0, implicit $mode, implicit $exec
+
+ ; MFMA_B: src2 = $vgpr4_vgpr5_vgpr6_vgpr7 (physical VGPR, declared live-in).
+ ; VISrc_128_f32 (src2 class of _vgprcd_e64) maps to VGPR_128 for physical
+ ; registers, so MIR verification accepts this operand.
+ ; The fix in initHeuristics detects !Src2->getReg().isVirtual() and skips
+ ; this MFMA, preventing the crash and the half-rewritten state.
+ %res_b:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 %10:av_128_align2, %11:av_128_align2, $vgpr4_vgpr5_vgpr6_vgpr7, 4, 4, %12.sub0:vreg_64_align2, %13:vgpr_32, 0, 0, implicit $mode, implicit $exec
+
+ ; Keep all regs live to maintain pressure and prevent DCE.
+ 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..0b41cf4877c8b
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-v7slice-pattern.ll
@@ -0,0 +1,278 @@
+; 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-COUNT-16: v_mfma_f32_16x16x32_f16 a[{{[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
+; BEFORE: bb.1.loop:
+; BEFORE: V_MFMA_F32_16X16X32_F16_vgprcd_e64
+; BEFORE-NOT: areg_128_align2
+; BEFORE: bb.2.epilogue:
+; BEFORE: V_MFMA_F32_16X16X32_F16_vgprcd_e64
+; BEFORE-NOT: V_MFMA_F32_16X16X32_F16_e64
+
+; stop-after=machine-scheduler: RewriteMFMAFormStage has fired.
+; All MFMAs promoted to _e64 (AGPR C/D), Case 3/2 COPYs inserted.
+; AFTER-LABEL: name: test_v7slice_scalar_phi_acc
+; AFTER: bb.0.entry:
+; Case 3: areg = COPY vreg (VGPR->AGPR, after non-MAI insertelement init).
+; AFTER: %{{[0-9]+}}:areg_128_align2 = COPY %{{[0-9]+}}
+; AFTER: bb.1.loop:
+; AFTER-NOT: V_MFMA_F32_16X16X32_F16_vgprcd_e64
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64
+; AFTER: bb.2.epilogue:
+; AFTER: %{{[0-9]+}}:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64
+; Case 2: vreg = COPY areg (AGPR->VGPR, before fptrunc non-MAI use in epilogue).
+; AFTER: %{{[0-9]+}}:vreg_128_align2 = COPY %{{[0-9]+}}
>From 8901416fa0ac08c0c84d748ae298e9ab306a74d2 Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Wed, 22 Apr 2026 07:21:32 +0000
Subject: [PATCH 2/2] [AMDGPU] Fix RewriteMFMAFormSchedStage
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 6 +++---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.h | 9 +++++++--
.../AMDGPU/rewrite-mfma-form-ReachingUseTracker.ll | 14 +++++++-------
3 files changed, 17 insertions(+), 12 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index 980beb38cd46d..0cd44b4404550 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -2251,7 +2251,7 @@ void GCNSchedStage::modifyRegionSchedule(unsigned RegionIdx,
}
void RewriteMFMAFormStage::resetRewriteCandsToVGPR(
- const std::vector<std::pair<MachineInstr *, unsigned>> &RewriteCands) {
+ ArrayRef<std::pair<MachineInstr *, unsigned>> RewriteCands) {
for (auto &[MI, OriginalOpcode] : RewriteCands) {
assert(TII->isMAI(*MI));
const TargetRegisterClass *ADefRC =
@@ -2371,7 +2371,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;
@@ -2473,7 +2473,7 @@ int64_t RewriteMFMAFormStage::getRewriteCost(
}
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;
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
index 258010a071704..09fe4f8f6a6b8 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
@@ -459,13 +459,18 @@ 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);
/// Resets all rewrite candidates in \p Cands back to their original VGPR
/// opcodes and register classes.
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-ReachingUseTracker.ll b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-ReachingUseTracker.ll
index 13bc16aad14a4..d613a18707c8d 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-ReachingUseTracker.ll
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-ReachingUseTracker.ll
@@ -1,10 +1,3 @@
-;
-; Verify the machine-scheduler-level rewrite() changes:
-; 1. All 9 MFMAs: opcode _vgprcd_e64 -> _e64, class vreg -> areg (AGPR form).
-; 2. Case 3 COPY: inserted after zeroinitializer def in loop.body (entry edge),
-; converts VGPR acc -> AGPR acc before MFMA (src2/dst reclassification).
-; 3. Case 2 COPY: inserted at exit block entry (ReachingUseTracker),
-; converts AGPR MFMA result -> VGPR before fptrunc (V_CVT_F16_F32).
; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a \
; RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
; RUN: < %s | FileCheck %s
@@ -13,6 +6,13 @@
; RUN: -stop-after=machine-scheduler \
; RUN: < %s | FileCheck %s --check-prefix=MIR
;
+; Verify the machine-scheduler-level rewrite() changes:
+; 1. All 9 MFMAs: opcode _vgprcd_e64 -> _e64, class vreg -> areg (AGPR form).
+; 2. Case 3 COPY: inserted after zeroinitializer def in loop.body (entry edge),
+; converts VGPR acc -> AGPR acc before MFMA (src2/dst reclassification).
+; 3. Case 2 COPY: inserted at exit block entry (ReachingUseTracker),
+; converts AGPR MFMA result -> VGPR before fptrunc (V_CVT_F16_F32).
+;
; CHECK-LABEL: test_ReachingUseTracker_crossblock_use:
; CHECK-COUNT-9: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], {{.*}}, {{.*}}, a[{{[0-9:]+}}]
;
More information about the llvm-commits
mailing list