[llvm] [AMDGPU] Fix RewriteMFMAFormSchedStage (PR #194887)
via llvm-commits
llvm-commits at lists.llvm.org
Fri May 1 00:01:18 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/4] [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/4] [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
>From bb383bab94c4ccd809d0ed4febb84542c8e65181 Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Thu, 30 Apr 2026 09:38:40 +0000
Subject: [PATCH 3/4] [AMDGPU] Fix resetRewriteCandsToVGPR subreg dest RC
corruption
When the MFMA dest operand has a subregister index (e.g. %13.sub0_sub1_sub2_sub3
on a vreg_1024_align2), TII->getRegClass(InstrDesc, 0) returns the constraint RC
for that subreg slot (vreg_128_align2), not the full virtual register's class.
This caused MRI.setRegClass to shrink the vreg RC, producing an invalid subreg
index assertion in MachineVerifier.
Fix by deriving VDefRC from the current AGPR class already stored in MRI via
getEquivalentVGPRClass, which correctly reflects the full register size.
Co-Authored-By: Claude Sonnet 4 (1M context) <noreply at anthropic.com>
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 5 +++--
1 file changed, 3 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index 0ad76ed9976a1..ff615bef0afed 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -2259,8 +2259,9 @@ 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);
+ 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));
>From 5dc6eb7a1f59c97518cf815f2119ccb6e48bd55f Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Fri, 1 May 2026 03:35:27 +0000
Subject: [PATCH 4/4] [AMDGPU] NFC: Revert comment spelling fixes in
GCNSchedStrategy
Restore comments in GCNSchedStrategy.{h,cpp} to their state before
the spelling corrections applied in 4e5a01ccb54e.
Co-Authored-By: Claude Sonnet 4 (1M context) <noreply at anthropic.com>
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 26 ++++++++++-----------
llvm/lib/Target/AMDGPU/GCNSchedStrategy.h | 17 +++++++-------
2 files changed, 21 insertions(+), 22 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index ff615bef0afed..2764779bceb77 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.
@@ -125,7 +125,7 @@ void GCNSchedStrategy::initialize(ScheduleDAGMI *DAG) {
Context->RegClassInfo->getNumAllocatableRegs(&AMDGPU::VGPR_32RegClass);
SIMachineFunctionInfo &MFI = *MF->getInfo<SIMachineFunctionInfo>();
- // Set the initial TargetOccupancy to the maximum occupancy that we can
+ // Set the initial TargetOccupnacy 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
@@ -340,7 +340,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.
@@ -350,7 +350,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.
@@ -371,7 +371,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.
@@ -849,8 +849,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.
@@ -1446,7 +1446,7 @@ Printable PreRARematStage::ScoredRemat::print() const {
bool PreRARematStage::initGCNSchedStage() {
// FIXME: This pass will invalidate cached BBLiveInMap and MBBLiveIns for
- // regions in between the defs and region we sunk the def to. Will need to be
+ // regions inbetween the defs and region we sinked the def to. Will need to be
// fixed if there is another pass after this pass.
assert(!S.hasNextStage());
@@ -1555,7 +1555,7 @@ bool PreRARematStage::initGCNSchedStage() {
}
// Rematerialize registers in successive rounds until all RP targets are
- // satisfied or until we run out of rematerialization candidates.
+ // satisifed or until we run out of rematerialization candidates.
BitVector RecomputeRP(DAG.Regions.size());
for (;;) {
RecomputeRP.reset();
@@ -1825,9 +1825,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 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.
+ // 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.
if (!DAG.RegionsWithHighRP[RegionIdx])
return false;
@@ -2391,7 +2391,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])
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
index 2ac958eef1cae..f128dcdbf1481 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
@@ -42,7 +42,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 {
@@ -105,7 +105,7 @@ class GCNSchedStrategy : public GenericScheduler {
// GCN RP Tracker for top-down scheduling
mutable GCNDownwardRPTracker DownwardTracker;
- // GCN RP Tracker for bottom-up scheduling
+ // GCN RP Tracker for botttom-up scheduling
mutable GCNUpwardRPTracker UpwardTracker;
bool UseGCNTrackers = false;
@@ -313,7 +313,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 retrieved with the regionIdx by calls to getLiveRegsForRegionIdx.
+ // be retreived with the regionIdx by calls to getLiveRegsForRegionIdx.
RegionPressureMap RegionLiveOuts;
// Return current region pressure.
@@ -393,7 +393,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,
@@ -446,7 +446,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,7 +464,6 @@ class RewriteMFMAFormStage : public GCNSchedStage {
/// 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 RewriteCands back to their original
/// VGPR opcodes and register classes.
void resetRewriteCandsToVGPR(
@@ -540,7 +539,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.
@@ -551,8 +550,8 @@ class PreRARematStage : public GCNSchedStage {
using RegisterIdx = Rematerializer::RegisterIdx;
/// A scored rematerialization candidate. Higher scores indicate more
- /// beneficial rematerializations. A null score indicates the
- /// rematerialization is not helpful to reduce RP in target regions.
+ /// beneficial rematerializations. A null score indicate the rematerialization
+ /// is not helpful to reduce RP in target regions.
struct ScoredRemat {
/// The register index handle in the rematerializer.
RegisterIdx RegIdx;
More information about the llvm-commits
mailing list