[llvm] [AMDGPU] Fix RewriteMFMAFormSchedStage (PR #194887)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Apr 29 08:34:07 PDT 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: xgxanq
<details>
<summary>Changes</summary>
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.
---
Patch is 39.37 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/194887.diff
4 Files Affected:
- (modified) llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp (+71-47)
- (modified) llvm/lib/Target/AMDGPU/GCNSchedStrategy.h (+6-2)
- (added) llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-check-half-rewrite.mir (+86)
- (added) llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-v7slice-pattern.ll (+380)
``````````diff
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...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/194887
More information about the llvm-commits
mailing list