[llvm] [AMDGPU] Fix RewriteMFMAFormSchedStage (PR #194887)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Apr 29 08:33:00 PDT 2026
https://github.com/xgxanq created https://github.com/llvm/llvm-project/pull/194887
1.RegionIdx is a field from the base class GCNSchedStage that tracks the current scheduling region index during normal scheduling. But rewrite() runs as a pre-pass (called from initGCNSchedStage(), before any region scheduling loop。
2.need to pre-check the dst and src2 in initHeuristics；
3.If the reaching use of the dst is a sub-reg, non-MFMA, and in the same basic block as the dst, we should also avoid creating redundant virtual registers and copies.
I can't create branches directly from llvm-project. As for stacked PRs, I can't implement that for now. So I forked two separate branches in my forked repo, and linked them together in the PR description. You can refer to https://github.com/llvm/llvm-project/pull/193429 to understand the fix just through the changes of the unit tests.
>From 6e08dcf6e65aa01150bda153fc09dbfaecd2b89b Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Wed, 29 Apr 2026 15:10:42 +0000
Subject: [PATCH] [AMDGPU] Fix RewriteMFMAFormSchedStage
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 118 +++---
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, 543 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..0cd44b4404550 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(
+ ArrayRef<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);
@@ -2339,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;
@@ -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,36 +2467,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 +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..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..d23ee6c7fec76
--- /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
+
+ %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.
+ %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) — triggers crash.
+ %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
+
+ 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]+}}
More information about the llvm-commits
mailing list