[llvm] [AMDGPU] Guard RewriteMFMAFormStage recolor against unsafe def/use (PR #217396)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Sep 14 02:47:07 PDT 2026
https://github.com/xgxanq updated https://github.com/llvm/llvm-project/pull/217396
>From 41b415b134f94032099b8ab24636287934e67b2c Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Wed, 19 Aug 2026 05:35:24 +0000
Subject: [PATCH 1/9] [AMDGPU] Guard RewriteMFMAFormStage recolor against
unsafe def/use
RewriteMFMAFormStage recolors an MFMA rewrite candidate's dst/src2 vreg to
AGPR, but previously only checked its uses. If the vreg is also redefined by a
non-MFMA VALU or feeds a VGPR-only use, the blind recolor creates an illegal
AGPR access and the machine verifier aborts.
Add isRecolorSafe(), which checks both the vreg's defs and their reaching uses
before recoloring. When the vreg is not safe to recolor, keep it in VGPR and
bridge it to AGPR with a copy for the MFMA operand instead.
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 210 ++++++++++-------
llvm/lib/Target/AMDGPU/GCNSchedStrategy.h | 15 +-
.../AMDGPU/mfma-no-multiple-copies.mir | 9 +-
.../AMDGPU/rewrite-mfma-form-fix-recolor.mir | 87 +++++++
.../AMDGPU/rewrite-mfma-single-exit-copy.ll | 219 +++++++++---------
.../AMDGPU/sched_mfma_rewrite_copies.mir | 105 +++++----
6 files changed, 393 insertions(+), 252 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-fix-recolor.mir
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index 3dd468772e9b1..59d4ae7ebadb2 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -2314,40 +2314,69 @@ void GCNSchedStage::modifyRegionSchedule(unsigned RegionIdx,
DAG.Regions[RegionIdx].first = MIOrder.front();
}
-/// Returns true if reaching def \p RD will be in AGPR form after the rewrite
-/// and so needs no bridge copy: a candidate MFMA in \p RewriteSet, an
-/// AV_MOV_*_IMM_PSEUDO, or a copy from a candidate src2 reg in \p CandSrc2Regs.
-/// A non-candidate MFMA stays in VGPR form and still needs a bridge.
-static bool isReachingDefAGPRForm(
- MachineInstr *RD, const SmallPtrSetImpl<MachineInstr *> &RewriteSet,
- const DenseSet<Register> &CandSrc2Regs, const SIInstrInfo &TII) {
- if (TII.isMAI(*RD))
- return RewriteSet.contains(RD);
- if (RD->getOpcode() == AMDGPU::AV_MOV_B32_IMM_PSEUDO ||
- RD->getOpcode() == AMDGPU::AV_MOV_B64_IMM_PSEUDO)
+static bool
+isRewriteCandidateMAI(const MachineInstr *MI, const SIInstrInfo *TII,
+ const SmallPtrSetImpl<MachineInstr *> &RewriteCandsSet) {
+ return TII->isMAI(*MI) && RewriteCandsSet.contains(MI);
+}
+
+static bool canWriteAGPR(const MachineInstr *MI, const SIInstrInfo *TII,
+ const SIRegisterInfo *SRI) {
+ // A writer can write its result into an AGPR lane iff its def operand's
+ // register-class constraint admits AGPRs (AV/agnostic classes do, plain VGPR
+ // classes don't). COPY is a generic opcode with no operand constraint, so
+ // special-case it.
+ if (MI->isCopy())
return true;
- if (RD->isCopy() && CandSrc2Regs.contains(RD->getOperand(1).getReg()))
- return true;
- return false;
+ const MCInstrDesc &Desc = MI->getDesc();
+ if (Desc.getNumDefs() == 0)
+ return false;
+ const TargetRegisterClass *DefRC = TII->getRegClass(Desc, 0);
+ return DefRC && SRI->hasAGPRs(DefRC);
}
-bool RewriteMFMAFormStage::hasUseRequiringVGPR(
- ArrayRef<SlotIndex> Src2ReachingDefs,
- const SmallPtrSetImpl<MachineInstr *> &RewriteSet) {
- for (SlotIndex RDIdx : Src2ReachingDefs) {
- const MachineInstr *RD = DAG.LIS->getInstructionFromIndex(RDIdx);
- SmallVector<MachineOperand *, 8> ReachingUses;
- findReachingUses(RD, DAG.LIS, ReachingUses);
- for (const MachineOperand *UseMO : ReachingUses) {
- const MachineInstr *UseMI = UseMO->getParent();
- if (UseMI->isCopy())
- continue;
- if (TII->isMAI(*UseMI) && RewriteSet.contains(UseMI))
+static bool useAcceptsAGPR(const MachineOperand *Use, const SIInstrInfo *TII,
+ const SIRegisterInfo *SRI) {
+ const MachineInstr *UseMI = Use->getParent();
+ // COPY is a generic opcode with no operand constraint; it can read an AGPR.
+ if (UseMI->isCopy())
+ return true;
+ // The use's static register-class constraint must admit AGPRs.
+ const TargetRegisterClass *UseRC =
+ UseMI->getRegClassConstraint(Use->getOperandNo(), TII, SRI);
+ return UseRC && SRI->hasAGPRs(UseRC);
+}
+
+bool RewriteMFMAFormStage::isRecolorSafe(
+ Register Reg, ArrayRef<MachineOperand *> DstReachingUses,
+ const SmallPtrSetImpl<MachineInstr *> &RewriteCandsSet, bool IsDst) {
+ for (MachineOperand &DefMO : DAG.MRI.def_operands(Reg)) {
+ MachineInstr *DefMI = DefMO.getParent();
+ // A candidate MFMA def is rewritten to AGPR form (it produces the AGPR
+ // result directly), so it does not constrain the recolor.
+ if (isRewriteCandidateMAI(DefMI, TII, RewriteCandsSet))
+ continue;
+ bool CanWriteAGPR = canWriteAGPR(DefMI, TII, SRI);
+ if (!CanWriteAGPR)
+ return false;
+ SmallVector<MachineOperand *, 8> DefReachingUses;
+ findReachingUses(DefMI, DAG.LIS, DefReachingUses);
+ for (MachineOperand *UseMO : DefReachingUses) {
+ // Exempt uses that do not constrain the recolor: for a dst, its own
+ // reaching uses (bridged to VGPR by case2); for a src2, uses that are
+ // candidate MFMAs (they read the AGPR result directly after rewrite).
+ if (IsDst) {
+ if (is_contained(DstReachingUses, UseMO))
+ continue;
+ } else if (isRewriteCandidateMAI(UseMO->getParent(), TII,
+ RewriteCandsSet)) {
continue;
- return true;
+ }
+ if (!useAcceptsAGPR(UseMO, TII, SRI))
+ return false;
}
}
- return false;
+ return true;
}
void RewriteMFMAFormStage::resetRewriteCandsToVGPR(
@@ -2378,12 +2407,6 @@ bool RewriteMFMAFormStage::isRewriteCandidate(MachineInstr *MI) const {
return false;
if (AMDGPU::getAGPRFormOp(MI->getOpcode()) == -1)
return false;
- // Reject candidates whose users force an unavoidable bridge copy.
- Register DstReg = MI->getOperand(0).getReg();
- for (const MachineInstr &UseMI : DAG.MRI.use_nodbg_instructions(DstReg)) {
- if (!TII->isMAI(UseMI) && !UseMI.isCopy())
- return false;
- }
return true;
}
@@ -2396,15 +2419,11 @@ bool RewriteMFMAFormStage::initHeuristics(
// Collect the candidate group, its members share AGPR-form operands
// post-rewrite, so reaching defs feeding any member don't need bridge copy.
SmallPtrSet<MachineInstr *, 16> RewriteSet;
- DenseSet<Register> CandSrc2Regs;
for (MachineBasicBlock &MBB : MF) {
for (MachineInstr &MI : MBB) {
if (!isRewriteCandidate(&MI))
continue;
RewriteSet.insert(&MI);
- MachineOperand *Src2 = TII->getNamedOperand(MI, AMDGPU::OpName::src2);
- if (Src2 && Src2->isReg())
- CandSrc2Regs.insert(Src2->getReg());
}
}
@@ -2425,17 +2444,13 @@ bool RewriteMFMAFormStage::initHeuristics(
SmallVector<SlotIndex, 8> Src2ReachingDefs;
findReachingDefs(*Src2, DAG.LIS, Src2ReachingDefs);
- // If src2 has a use that must remain VGPR, it cannot be reclassified to
- // AGPR.
- bool Src2NeedsVGPR = hasUseRequiringVGPR(Src2ReachingDefs, RewriteSet);
- Src2NeedsVGPRCache[&MI] = Src2NeedsVGPR;
-
- for (SlotIndex RDIdx : Src2ReachingDefs) {
- MachineInstr *RD = DAG.LIS->getInstructionFromIndex(RDIdx);
- if (!Src2NeedsVGPR &&
- isReachingDefAGPRForm(RD, RewriteSet, CandSrc2Regs, *TII))
- continue;
- CopyForDef.insert(RD);
+ bool Src2RecolorSafe =
+ isRecolorSafe(Src2->getReg(), {}, RewriteSet, /*IsDst=*/false);
+ if (!Src2RecolorSafe) {
+ for (SlotIndex RDIdx : Src2ReachingDefs) {
+ MachineInstr *RD = DAG.LIS->getInstructionFromIndex(RDIdx);
+ CopyForDef.insert(RD);
+ }
}
}
@@ -2443,33 +2458,39 @@ bool RewriteMFMAFormStage::initHeuristics(
SmallVector<MachineOperand *, 8> DstReachingUses;
findReachingUses(&MI, DAG.LIS, DstReachingUses);
+ bool DstRecolorSafe =
+ isRecolorSafe(Dst.getReg(), DstReachingUses, RewriteSet,
+ /*IsDst=*/true);
for (MachineOperand *RUOp : DstReachingUses) {
MachineInstr *UserMI = RUOp->getParent();
+ bool NeedsAGPRToVGPRCopy = true;
// Group members read the AGPR result directly.
if (TII->isMAI(*UserMI) && RewriteSet.contains(UserMI))
- continue;
+ NeedsAGPRToVGPRCopy = false;
// For any user of the result of the MFMA which is not an MFMA, we
- // insert a copy. For a given register, we will only insert one copy
- // per user block.
- CopyForUse[UserMI->getParent()].insert(RUOp->getReg());
-
- if (TII->isMAI(*UserMI))
+ // record a copy location. For a given register, we only record one
+ // copy per user block.
+ if (NeedsAGPRToVGPRCopy)
+ CopyForUse[UserMI->getParent()].insert(RUOp->getReg());
+
+ // If the dst can be wholly recolored to AGPR, its reaching defs are
+ // reclassified along with it, so no per-def bridge copy is needed.
+ if (DstRecolorSafe)
continue;
-
SmallVector<SlotIndex, 8> DstUsesReachingDefs;
findReachingDefs(*RUOp, DAG.LIS, DstUsesReachingDefs);
for (SlotIndex RDIndex : DstUsesReachingDefs) {
MachineInstr *RD = DAG.LIS->getInstructionFromIndex(RDIndex);
- if (TII->isMAI(*RD))
+ if (TII->isMAI(*RD) && RewriteSet.contains(RD))
continue;
- // For any definition of the user of the MFMA which is not an MFMA,
- // we insert a copy. We do this to transform all the reaching defs
- // of this use to AGPR. By doing this, we can insert a copy from
- // AGPR to VGPR at the user rather than after the MFMA.
+ // Dst cannot be recolored: for any non-MFMA reaching def of this
+ // use, record a copy location. This transforms all the reaching
+ // defs of this use to AGPR so the AGPR-to-VGPR copy lands at the
+ // user rather than after the MFMA.
CopyForDef.insert(RD);
}
}
@@ -2691,17 +2712,21 @@ bool RewriteMFMAFormStage::rewrite(
findReachingDefs(*Src2, DAG.LIS, Src2ReachingDefs);
SmallSetVector<MachineInstr *, 8> Src2DefsReplace;
- // If src2 has a use that must remain VGPR, it cannot be reclassified to
- // AGPR.
- bool Src2NeedsVGPR = Src2NeedsVGPRCache.lookup(MI);
-
- for (SlotIndex RDIndex : Src2ReachingDefs) {
- MachineInstr *RD = DAG.LIS->getInstructionFromIndex(RDIndex);
- if (!Src2NeedsVGPR &&
- isReachingDefAGPRForm(RD, RewriteCandsSet, RewriteSrc2Regs, *TII))
- continue;
-
- Src2DefsReplace.insert(RD);
+ // An already-redefined src2 reuses its mapped reg below, so treat it as
+ // unsafe to recolor and collect the reaching defs that need a bridge.
+ bool Src2AlreadyRedef = RedefMap.contains(Src2Reg);
+ bool Src2RecolorSafe =
+ !Src2AlreadyRedef &&
+ isRecolorSafe(Src2Reg, {}, RewriteCandsSet, /*IsDst=*/false);
+ // src2 cannot be recolored to AGPR: collect its non-candidate reaching
+ // defs to bridge below (candidate MFMAs already produce AGPR directly).
+ if (!Src2RecolorSafe) {
+ for (SlotIndex RDIndex : Src2ReachingDefs) {
+ MachineInstr *RD = DAG.LIS->getInstructionFromIndex(RDIndex);
+ if (isRewriteCandidateMAI(RD, TII, RewriteCandsSet))
+ continue;
+ Src2DefsReplace.insert(RD);
+ }
}
if (!Src2DefsReplace.empty()) {
@@ -2766,34 +2791,47 @@ bool RewriteMFMAFormStage::rewrite(
findReachingUses(MI, DAG.LIS, DstReachingUses);
+ // An already-redefined dst reuses its mapped reg, so treat it as unsafe to
+ // recolor and bridge its reaching defs instead.
+ bool DstAlreadyRedef = RedefMap.contains(DstReg);
+ bool DstRecolorSafe =
+ !DstAlreadyRedef &&
+ isRecolorSafe(DstReg, DstReachingUses, RewriteCandsSet, /*IsDst=*/true);
for (MachineOperand *RUOp : DstReachingUses) {
MachineInstr *UserMI = RUOp->getParent();
- // Group members read the AGPR result directly.
- if (TII->isMAI(*UserMI) && RewriteCandsSet.contains(UserMI))
- continue;
-
- // If there is a non mai reaching use, then we need a copy.
- if (find(DstReachingUseCopies, RUOp) == DstReachingUseCopies.end())
+ // A group-member MFMA reads the AGPR result directly; any other user can
+ // only avoid a bridge if it accepts an AGPR operand.
+ bool CanReadAGPR = TII->isMAI(*UserMI) ? RewriteCandsSet.contains(UserMI)
+ : useAcceptsAGPR(RUOp, TII, SRI);
+ if (!CanReadAGPR &&
+ find(DstReachingUseCopies, RUOp) == DstReachingUseCopies.end())
DstReachingUseCopies.push_back(RUOp);
-
- // Non-rewritten MAI: its defs aren't being reclassified.
- if (TII->isMAI(*UserMI))
+ // If the dst is wholly recolored to AGPR, its reaching defs are
+ // reclassified along with it, so none of them need a bridge copy.
+ if (DstRecolorSafe)
continue;
-
SmallVector<SlotIndex, 8> DstUsesReachingDefs;
findReachingDefs(*RUOp, DAG.LIS, DstUsesReachingDefs);
for (SlotIndex RDIndex : DstUsesReachingDefs) {
MachineInstr *RD = DAG.LIS->getInstructionFromIndex(RDIndex);
- if (TII->isMAI(*RD))
+ if (isRewriteCandidateMAI(RD, TII, RewriteCandsSet))
continue;
-
- // If there is a non mai reaching def of this reaching use, then we will
- // need a copy.
+ // A non-candidate reaching def must be bridged to VGPR; record it once
+ // (dedup against DstUseDefsReplace).
if (find(DstUseDefsReplace, RD) == DstUseDefsReplace.end())
DstUseDefsReplace.push_back(RD);
}
}
+ // The dst has no reaching uses and cannot be recolored: create a fresh
+ // reg to carry the AGPR-form value and record the mapping, leaving the
+ // original dst reg in VGPR form.
+ if (DstReachingUses.empty() && !DstRecolorSafe) {
+ const TargetRegisterClass *DstRC = DAG.MRI.getRegClass(DstReg);
+ const TargetRegisterClass *VGPRRC = SRI->getEquivalentVGPRClass(DstRC);
+ MappedReg = DAG.MRI.createVirtualRegister(VGPRRC);
+ RedefMap[DstReg] = MappedReg;
+ }
if (!DstUseDefsReplace.empty()) {
auto RI = RedefMap.find(DstReg);
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
index 2059f4e6479ff..f922b694dd6c2 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
@@ -443,10 +443,6 @@ class RewriteMFMAFormStage : public GCNSchedStage {
const SIInstrInfo *TII;
const SIRegisterInfo *SRI;
- /// Per-candidate cache of the src2 "needs VGPR" decision, computed once
- /// and reused on-demand.
- DenseMap<const MachineInstr *, bool> Src2NeedsVGPRCache;
-
/// 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
@@ -486,11 +482,12 @@ class RewriteMFMAFormStage : public GCNSchedStage {
void findReachingUses(const MachineInstr *DefMI, LiveIntervals *LIS,
SmallVectorImpl<MachineOperand *> &ReachingUses);
- /// Returns true if the src2 register with reaching defs \p Src2ReachingDefs
- /// has a use other than a group MFMA (in \p RewriteSet) or a copy, which
- /// would keep it in VGPR form rather than let it be reclassified to AGPR.
- bool hasUseRequiringVGPR(ArrayRef<SlotIndex> Src2ReachingDefs,
- const SmallPtrSetImpl<MachineInstr *> &RewriteSet);
+ /// Returns true if \p Reg (a rewrite candidate's dst or src2, selected by
+ /// \p IsDst) can be wholly reclassified to AGPR without creating an illegal
+ /// VALU AGPR access.
+ bool isRecolorSafe(Register Reg, ArrayRef<MachineOperand *> DstReachingUses,
+ const SmallPtrSetImpl<MachineInstr *> &RewriteCandsSet,
+ bool IsDst);
public:
bool initGCNSchedStage() override;
diff --git a/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir b/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir
index c7ae0a2c2ef93..de9c11753db37 100644
--- a/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir
+++ b/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir
@@ -42,10 +42,11 @@ body: |
; CHECK-NEXT: bb.2:
; CHECK-NEXT: KILL %p0, %p1, %p2, %p3, %p4, %p5, %p6, %p7
; CHECK-NEXT: %m4:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m2, 0, 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY %m4
- ; CHECK-NEXT: %out1:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, [[COPY1]], 0, 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: %out2:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s1:vreg_128_align2, undef %s0:vreg_128_align2, [[COPY1]], 0, 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: S_ENDPGM 0, implicit %out1, implicit %out2
+ ; CHECK-NEXT: %out1:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m4, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %out2:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s1:vreg_128_align2, undef %s0:vreg_128_align2, %m4, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY %out1
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:vreg_128_align2 = COPY %out2
+ ; CHECK-NEXT: S_ENDPGM 0, implicit [[COPY1]], implicit [[COPY2]]
bb.0:
successors: %bb.1(0x80000000)
%p0:vreg_1024 = IMPLICIT_DEF
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-fix-recolor.mir b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-fix-recolor.mir
new file mode 100644
index 0000000000000..cc53c87fd6299
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-fix-recolor.mir
@@ -0,0 +1,87 @@
+# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 6
+# RUN: llc -mtriple=amdgpu9.50-amd-amdhsa -run-pass=machine-scheduler \
+# RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+# RUN: -verify-machineinstrs -o - %s | FileCheck %s
+
+# Correctness test for the RewriteMFMAFormStage def-side recolor guard. The loop
+# MFMA's src2 %acc is partially redefined in the epilogue (%acc.sub0) and read by
+# a VGPR-only DS_READ_B32. isRecolorSafe scans the reaching uses of %acc's defs,
+# sees that read, and keeps %acc in VGPR -- bridging it to AGPR with a copy for
+# src2 instead of recoloring the whole vreg (which would abort with
+# "AReg_64_Align2.sub0 cannot be used for VGPR_32 operands").
+
+--- |
+ define amdgpu_kernel void @mfma_fix_recolor() #0 {
+ ret void
+ }
+
+ attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }
+...
+---
+name: mfma_fix_recolor
+tracksRegLiveness: true
+body: |
+ ; CHECK-LABEL: name: mfma_fix_recolor
+ ; CHECK: bb.0:
+ ; CHECK-NEXT: successors: %bb.1(0x80000000)
+ ; CHECK-NEXT: liveins: $vgpr0_vgpr1, $vgpr2_vgpr3, $sgpr0
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: %s0:vreg_64_align2 = COPY $vgpr0_vgpr1
+ ; CHECK-NEXT: %s1:vreg_64_align2 = COPY $vgpr2_vgpr3
+ ; CHECK-NEXT: %cnt:sgpr_32 = COPY $sgpr0
+ ; CHECK-NEXT: %acc:vreg_64_align2 = AV_MOV_B64_IMM_PSEUDO 0, implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_64_align2 = COPY %acc
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: bb.1:
+ ; CHECK-NEXT: successors: %bb.1(0x40000000), %bb.2(0x40000000)
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: dead %mdst:areg_64_align2 = V_MFMA_F64_4X4X4F64_e64 %s0, %s1, [[COPY]], 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %cnt:sgpr_32 = S_SUB_I32 %cnt, 1, implicit-def $scc
+ ; CHECK-NEXT: S_CBRANCH_SCC1 %bb.1, implicit $scc
+ ; CHECK-NEXT: S_BRANCH %bb.2
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: bb.2:
+ ; CHECK-NEXT: %acc.sub0:vreg_64_align2 = AV_MOV_B32_IMM_PSEUDO 7, implicit $exec
+ ; CHECK-NEXT: dead %conn:vreg_64_align2 = COPY %acc
+ ; CHECK-NEXT: %ld:vgpr_32 = DS_READ_B32_gfx9 %acc.sub0, 0, 0, implicit $exec :: (load (s32), addrspace 3)
+ ; CHECK-NEXT: [[DEF:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF1:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF2:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF3:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF4:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF5:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF6:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF7:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF8:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: S_ENDPGM 0, implicit %ld, implicit [[DEF]], implicit [[DEF1]], implicit [[DEF2]], implicit [[DEF3]], implicit [[DEF4]], implicit [[DEF5]], implicit [[DEF6]], implicit [[DEF7]], implicit [[DEF8]]
+ bb.0:
+ successors: %bb.1
+ liveins: $vgpr0_vgpr1, $vgpr2_vgpr3, $sgpr0
+
+ %s0:vreg_64_align2 = COPY $vgpr0_vgpr1
+ %s1:vreg_64_align2 = COPY $vgpr2_vgpr3
+ %cnt:sgpr_32 = COPY $sgpr0
+ %p0:vreg_1024 = IMPLICIT_DEF
+ %p1:vreg_1024 = IMPLICIT_DEF
+ %p2:vreg_1024 = IMPLICIT_DEF
+ %p3:vreg_1024 = IMPLICIT_DEF
+ %p4:vreg_1024 = IMPLICIT_DEF
+ %p5:vreg_1024 = IMPLICIT_DEF
+ %p6:vreg_1024 = IMPLICIT_DEF
+ %p7:vreg_1024 = IMPLICIT_DEF
+ %p8:vreg_1024 = IMPLICIT_DEF
+ %acc:vreg_64_align2 = AV_MOV_B64_IMM_PSEUDO 0, implicit $exec
+
+ bb.1:
+ successors: %bb.1, %bb.2
+ dead %mdst:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %s0, %s1, %acc, 0, 0, 0, implicit $mode, implicit $exec
+ %cnt:sgpr_32 = S_SUB_I32 %cnt, 1, implicit-def $scc
+ S_CBRANCH_SCC1 %bb.1, implicit $scc
+ S_BRANCH %bb.2
+
+ bb.2:
+ %acc.sub0 = AV_MOV_B32_IMM_PSEUDO 7, implicit $exec
+ %ld:vgpr_32 = DS_READ_B32_gfx9 %acc.sub0, 0, 0, implicit $exec :: (load (s32), addrspace 3)
+ dead %conn:vreg_64_align2 = COPY %acc
+ S_ENDPGM 0, implicit %ld, implicit %p0, implicit %p1, implicit %p2, implicit %p3, implicit %p4, implicit %p5, implicit %p6, implicit %p7, implicit %p8
+...
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-single-exit-copy.ll b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-single-exit-copy.ll
index cfbb80c7696a0..b0065a6fc750d 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-single-exit-copy.ll
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-single-exit-copy.ll
@@ -250,14 +250,6 @@ define amdgpu_kernel void @single_exit_copy(
; CHECK-NEXT: v_mov_b32_e32 v194, 0
; CHECK-NEXT: v_mov_b32_e32 v193, 0
; CHECK-NEXT: v_mov_b32_e32 v192, 0
-; CHECK-NEXT: v_mov_b32_e32 v243, 0
-; CHECK-NEXT: v_mov_b32_e32 v242, 0
-; CHECK-NEXT: v_mov_b32_e32 v241, 0
-; CHECK-NEXT: v_mov_b32_e32 v240, 0
-; CHECK-NEXT: v_mov_b32_e32 v247, 0
-; CHECK-NEXT: v_mov_b32_e32 v246, 0
-; CHECK-NEXT: v_mov_b32_e32 v245, 0
-; CHECK-NEXT: v_mov_b32_e32 v244, 0
; CHECK-NEXT: v_mov_b32_e32 v251, 0
; CHECK-NEXT: v_mov_b32_e32 v250, 0
; CHECK-NEXT: v_mov_b32_e32 v249, 0
@@ -266,6 +258,14 @@ define amdgpu_kernel void @single_exit_copy(
; CHECK-NEXT: v_mov_b32_e32 v254, 0
; CHECK-NEXT: v_mov_b32_e32 v253, 0
; CHECK-NEXT: v_mov_b32_e32 v252, 0
+; CHECK-NEXT: v_mov_b32_e32 v243, 0
+; CHECK-NEXT: v_mov_b32_e32 v242, 0
+; CHECK-NEXT: v_mov_b32_e32 v241, 0
+; CHECK-NEXT: v_mov_b32_e32 v240, 0
+; CHECK-NEXT: v_mov_b32_e32 v247, 0
+; CHECK-NEXT: v_mov_b32_e32 v246, 0
+; CHECK-NEXT: v_mov_b32_e32 v245, 0
+; CHECK-NEXT: v_mov_b32_e32 v244, 0
; CHECK-NEXT: v_mov_b32_e32 v239, 0
; CHECK-NEXT: v_mov_b32_e32 v238, 0
; CHECK-NEXT: v_mov_b32_e32 v237, 0
@@ -304,18 +304,18 @@ define amdgpu_kernel void @single_exit_copy(
; CHECK-NEXT: s_nop 1
; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], a[4:7], a[8:11], a[0:3]
; CHECK-NEXT: s_add_i32 s1, s1, 1
-; CHECK-NEXT: v_pk_add_f32 v[242:243], v[242:243], v[242:243]
-; CHECK-NEXT: v_pk_add_f32 v[240:241], v[240:241], v[240:241]
-; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], a[4:7], a[8:11], a[0:3]
-; CHECK-NEXT: v_add_f32_e64 v246, v246, v246
-; CHECK-NEXT: v_add_f32_e64 v247, v247, v247
-; CHECK-NEXT: v_pk_add_f32 v[244:245], v[244:245], v[244:245]
; CHECK-NEXT: v_pk_add_f32 v[250:251], v[250:251], v[250:251]
+; CHECK-NEXT: v_pk_add_f32 v[248:249], v[248:249], v[248:249]
; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], a[4:7], a[8:11], a[0:3]
-; CHECK-NEXT: v_add_f32_e64 v248, v248, v248
-; CHECK-NEXT: v_add_f32_e64 v249, v249, v249
-; CHECK-NEXT: v_pk_add_f32 v[254:255], v[254:255], v[254:255]
+; CHECK-NEXT: v_add_f32_e64 v254, v254, v254
+; CHECK-NEXT: v_add_f32_e64 v255, v255, v255
; CHECK-NEXT: v_pk_add_f32 v[252:253], v[252:253], v[252:253]
+; CHECK-NEXT: v_pk_add_f32 v[242:243], v[242:243], v[242:243]
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], a[4:7], a[8:11], a[0:3]
+; CHECK-NEXT: v_add_f32_e64 v240, v240, v240
+; CHECK-NEXT: v_add_f32_e64 v241, v241, v241
+; CHECK-NEXT: v_pk_add_f32 v[246:247], v[246:247], v[246:247]
+; CHECK-NEXT: v_pk_add_f32 v[244:245], v[244:245], v[244:245]
; CHECK-NEXT: v_pk_add_f32 v[238:239], v[238:239], v[238:239]
; CHECK-NEXT: v_pk_add_f32 v[236:237], v[236:237], v[236:237]
; CHECK-NEXT: v_pk_add_f32 v[234:235], v[234:235], v[234:235]
@@ -441,104 +441,105 @@ define amdgpu_kernel void @single_exit_copy(
; CHECK-NEXT: ; %bb.2: ; %exit
; CHECK-NEXT: .Ltmp1:
; CHECK-NEXT: .loc 1 10 1 is_stmt 1 ; test:10:1
-; CHECK-NEXT: v_accvgpr_write_b32 a4, s8
-; CHECK-NEXT: v_accvgpr_write_b32 a5, s9
-; CHECK-NEXT: v_accvgpr_write_b32 a6, s10
-; CHECK-NEXT: v_accvgpr_write_b32 a7, s11
+; CHECK-NEXT: v_accvgpr_write_b32 a8, s8
+; CHECK-NEXT: v_accvgpr_write_b32 a9, s9
+; CHECK-NEXT: v_accvgpr_write_b32 a10, s10
+; CHECK-NEXT: v_accvgpr_write_b32 a11, s11
+; CHECK-NEXT: v_accvgpr_write_b32 a12, s12
+; CHECK-NEXT: v_accvgpr_write_b32 a13, s13
+; CHECK-NEXT: v_accvgpr_write_b32 a14, s14
+; CHECK-NEXT: v_accvgpr_write_b32 a15, s15
; CHECK-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x28
-; CHECK-NEXT: v_accvgpr_write_b32 a8, s12
-; CHECK-NEXT: v_accvgpr_write_b32 a9, s13
-; CHECK-NEXT: v_accvgpr_write_b32 a10, s14
-; CHECK-NEXT: v_accvgpr_write_b32 a11, s15
-; CHECK-NEXT: v_accvgpr_write_b32 a15, v3
-; CHECK-NEXT: v_accvgpr_write_b32 a14, v2
-; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], a[4:7], a[8:11], a[0:3]
-; CHECK-NEXT: v_accvgpr_write_b32 a13, v1
-; CHECK-NEXT: v_accvgpr_write_b32 a12, v0
+; CHECK-NEXT: v_accvgpr_write_b32 a19, v3
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], a[8:11], a[12:15], a[0:3]
+; CHECK-NEXT: v_accvgpr_write_b32 a18, v2
+; CHECK-NEXT: v_accvgpr_write_b32 a17, v1
+; CHECK-NEXT: v_accvgpr_write_b32 a16, v0
+; CHECK-NEXT: .loc 1 20 1 ; test:20:1
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[4:7], a[8:11], a[12:15], a[0:3]
; CHECK-NEXT: v_mov_b32_e32 v0, 0
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
-; CHECK-NEXT: global_store_dwordx4 v0, v[240:243], s[0:1] offset:2032
-; CHECK-NEXT: v_mov_b32_e32 v0, 0
-; CHECK-NEXT: global_store_dwordx4 v0, v[244:247], s[0:1] offset:2016
-; CHECK-NEXT: global_store_dwordx4 v0, v[248:251], s[0:1] offset:2000
+; CHECK-NEXT: global_store_dwordx4 v0, v[248:251], s[0:1] offset:2032
; CHECK-NEXT: v_mov_b32_e32 v0, 0
-; CHECK-NEXT: v_accvgpr_read_b32 v243, a3
-; CHECK-NEXT: v_accvgpr_read_b32 v242, a2
-; CHECK-NEXT: v_accvgpr_read_b32 v241, a1
-; CHECK-NEXT: v_accvgpr_read_b32 v240, a0
-; CHECK-NEXT: global_store_dwordx4 v0, v[252:255], s[0:1] offset:1984
-; CHECK-NEXT: .loc 1 20 1 ; test:20:1
-; CHECK-NEXT: s_nop 0
-; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[244:247], a[4:7], a[8:11], v[240:243]
-; CHECK-NEXT: s_nop 7
-; CHECK-NEXT: v_mov_b32_e32 v245, 0
-; CHECK-NEXT: global_store_dwordx4 v245, v[236:239], s[0:1] offset:1968
; CHECK-NEXT: .loc 1 30 1 ; test:30:1
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], a[12:15], a[8:11], a[0:3]
+; CHECK-NEXT: v_mov_b32_e32 v248, 0
+; CHECK-NEXT: global_store_dwordx4 v248, v[240:243], s[0:1] offset:2000
+; CHECK-NEXT: global_store_dwordx4 v248, v[244:247], s[0:1] offset:1984
+; CHECK-NEXT: global_store_dwordx4 v0, v[252:255], s[0:1] offset:2016
+; CHECK-NEXT: v_mov_b32_e32 v242, 0
+; CHECK-NEXT: global_store_dwordx4 v242, v[236:239], s[0:1] offset:1968
; CHECK-NEXT: s_nop 1
-; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[236:239], a[8:11], a[4:7], v[240:243]
-; CHECK-NEXT: s_nop 7
-; CHECK-NEXT: v_add_f32_e32 v238, 1.0, v244
-; CHECK-NEXT: v_add_f32_e32 v239, 2.0, v236
-; CHECK-NEXT: global_store_dwordx2 v245, v[238:239], s[0:1]
-; CHECK-NEXT: global_store_dwordx4 v245, v[232:235], s[0:1] offset:1952
-; CHECK-NEXT: global_store_dwordx4 v245, v[228:231], s[0:1] offset:1936
-; CHECK-NEXT: global_store_dwordx4 v245, v[224:227], s[0:1] offset:1920
-; CHECK-NEXT: global_store_dwordx4 v245, v[220:223], s[0:1] offset:368
-; CHECK-NEXT: global_store_dwordx4 v245, v[216:219], s[0:1] offset:352
-; CHECK-NEXT: global_store_dwordx4 v245, v[212:215], s[0:1] offset:336
-; CHECK-NEXT: global_store_dwordx4 v245, v[208:211], s[0:1] offset:320
-; CHECK-NEXT: global_store_dwordx4 v245, v[204:207], s[0:1] offset:304
-; CHECK-NEXT: global_store_dwordx4 v245, v[200:203], s[0:1] offset:288
-; CHECK-NEXT: global_store_dwordx4 v245, v[196:199], s[0:1] offset:272
-; CHECK-NEXT: global_store_dwordx4 v245, v[192:195], s[0:1] offset:256
-; CHECK-NEXT: global_store_dwordx4 v245, v[188:191], s[0:1] offset:496
-; CHECK-NEXT: global_store_dwordx4 v245, v[184:187], s[0:1] offset:480
-; CHECK-NEXT: global_store_dwordx4 v245, v[180:183], s[0:1] offset:464
-; CHECK-NEXT: global_store_dwordx4 v245, v[176:179], s[0:1] offset:448
-; CHECK-NEXT: global_store_dwordx4 v245, v[172:175], s[0:1] offset:432
-; CHECK-NEXT: global_store_dwordx4 v245, v[168:171], s[0:1] offset:416
-; CHECK-NEXT: global_store_dwordx4 v245, v[164:167], s[0:1] offset:400
-; CHECK-NEXT: global_store_dwordx4 v245, v[160:163], s[0:1] offset:384
-; CHECK-NEXT: global_store_dwordx4 v245, v[156:159], s[0:1] offset:624
-; CHECK-NEXT: global_store_dwordx4 v245, v[152:155], s[0:1] offset:608
-; CHECK-NEXT: global_store_dwordx4 v245, v[148:151], s[0:1] offset:592
-; CHECK-NEXT: global_store_dwordx4 v245, v[144:147], s[0:1] offset:576
-; CHECK-NEXT: global_store_dwordx4 v245, v[140:143], s[0:1] offset:560
-; CHECK-NEXT: global_store_dwordx4 v245, v[136:139], s[0:1] offset:544
-; CHECK-NEXT: global_store_dwordx4 v245, v[132:135], s[0:1] offset:528
-; CHECK-NEXT: global_store_dwordx4 v245, v[128:131], s[0:1] offset:512
-; CHECK-NEXT: global_store_dwordx4 v245, v[124:127], s[0:1] offset:752
-; CHECK-NEXT: global_store_dwordx4 v245, v[120:123], s[0:1] offset:736
-; CHECK-NEXT: global_store_dwordx4 v245, v[116:119], s[0:1] offset:720
-; CHECK-NEXT: global_store_dwordx4 v245, v[112:115], s[0:1] offset:704
-; CHECK-NEXT: global_store_dwordx4 v245, v[108:111], s[0:1] offset:688
-; CHECK-NEXT: global_store_dwordx4 v245, v[104:107], s[0:1] offset:672
-; CHECK-NEXT: global_store_dwordx4 v245, v[100:103], s[0:1] offset:656
-; CHECK-NEXT: global_store_dwordx4 v245, v[96:99], s[0:1] offset:640
-; CHECK-NEXT: global_store_dwordx4 v245, v[92:95], s[0:1] offset:880
-; CHECK-NEXT: global_store_dwordx4 v245, v[88:91], s[0:1] offset:864
-; CHECK-NEXT: global_store_dwordx4 v245, v[84:87], s[0:1] offset:848
-; CHECK-NEXT: global_store_dwordx4 v245, v[80:83], s[0:1] offset:832
-; CHECK-NEXT: global_store_dwordx4 v245, v[76:79], s[0:1] offset:816
-; CHECK-NEXT: global_store_dwordx4 v245, v[72:75], s[0:1] offset:800
-; CHECK-NEXT: global_store_dwordx4 v245, v[68:71], s[0:1] offset:784
-; CHECK-NEXT: global_store_dwordx4 v245, v[64:67], s[0:1] offset:768
-; CHECK-NEXT: global_store_dwordx4 v245, v[60:63], s[0:1] offset:1008
-; CHECK-NEXT: global_store_dwordx4 v245, v[56:59], s[0:1] offset:992
-; CHECK-NEXT: global_store_dwordx4 v245, v[52:55], s[0:1] offset:976
-; CHECK-NEXT: global_store_dwordx4 v245, v[48:51], s[0:1] offset:960
-; CHECK-NEXT: global_store_dwordx4 v245, v[44:47], s[0:1] offset:944
-; CHECK-NEXT: global_store_dwordx4 v245, v[40:43], s[0:1] offset:928
-; CHECK-NEXT: global_store_dwordx4 v245, v[36:39], s[0:1] offset:912
-; CHECK-NEXT: global_store_dwordx4 v245, v[32:35], s[0:1] offset:896
-; CHECK-NEXT: global_store_dwordx4 v245, v[28:31], s[0:1] offset:1136
-; CHECK-NEXT: global_store_dwordx4 v245, v[24:27], s[0:1] offset:1120
-; CHECK-NEXT: global_store_dwordx4 v245, v[20:23], s[0:1] offset:1104
-; CHECK-NEXT: global_store_dwordx4 v245, v[16:19], s[0:1] offset:1088
-; CHECK-NEXT: global_store_dwordx4 v245, v[12:15], s[0:1] offset:1072
-; CHECK-NEXT: global_store_dwordx4 v245, v[8:11], s[0:1] offset:1056
-; CHECK-NEXT: global_store_dwordx4 v245, v[4:7], s[0:1] offset:1040
-; CHECK-NEXT: global_store_dwordx4 v245, a[12:15], s[0:1] offset:1024
+; CHECK-NEXT: v_accvgpr_read_b32 v238, a6
+; CHECK-NEXT: v_accvgpr_read_b32 v237, a5
+; CHECK-NEXT: v_accvgpr_read_b32 v236, a4
+; CHECK-NEXT: v_accvgpr_read_b32 v238, a0
+; CHECK-NEXT: v_accvgpr_read_b32 v239, a7
+; CHECK-NEXT: v_add_f32_e32 v236, 1.0, v236
+; CHECK-NEXT: v_add_f32_e32 v237, 2.0, v238
+; CHECK-NEXT: v_accvgpr_read_b32 v241, a3
+; CHECK-NEXT: v_accvgpr_read_b32 v240, a2
+; CHECK-NEXT: v_accvgpr_read_b32 v239, a1
+; CHECK-NEXT: global_store_dwordx2 v242, v[236:237], s[0:1]
+; CHECK-NEXT: global_store_dwordx4 v242, v[232:235], s[0:1] offset:1952
+; CHECK-NEXT: global_store_dwordx4 v242, v[228:231], s[0:1] offset:1936
+; CHECK-NEXT: global_store_dwordx4 v242, v[224:227], s[0:1] offset:1920
+; CHECK-NEXT: global_store_dwordx4 v242, v[220:223], s[0:1] offset:368
+; CHECK-NEXT: global_store_dwordx4 v242, v[216:219], s[0:1] offset:352
+; CHECK-NEXT: global_store_dwordx4 v242, v[212:215], s[0:1] offset:336
+; CHECK-NEXT: global_store_dwordx4 v242, v[208:211], s[0:1] offset:320
+; CHECK-NEXT: global_store_dwordx4 v242, v[204:207], s[0:1] offset:304
+; CHECK-NEXT: global_store_dwordx4 v242, v[200:203], s[0:1] offset:288
+; CHECK-NEXT: global_store_dwordx4 v242, v[196:199], s[0:1] offset:272
+; CHECK-NEXT: global_store_dwordx4 v242, v[192:195], s[0:1] offset:256
+; CHECK-NEXT: global_store_dwordx4 v242, v[188:191], s[0:1] offset:496
+; CHECK-NEXT: global_store_dwordx4 v242, v[184:187], s[0:1] offset:480
+; CHECK-NEXT: global_store_dwordx4 v242, v[180:183], s[0:1] offset:464
+; CHECK-NEXT: global_store_dwordx4 v242, v[176:179], s[0:1] offset:448
+; CHECK-NEXT: global_store_dwordx4 v242, v[172:175], s[0:1] offset:432
+; CHECK-NEXT: global_store_dwordx4 v242, v[168:171], s[0:1] offset:416
+; CHECK-NEXT: global_store_dwordx4 v242, v[164:167], s[0:1] offset:400
+; CHECK-NEXT: global_store_dwordx4 v242, v[160:163], s[0:1] offset:384
+; CHECK-NEXT: global_store_dwordx4 v242, v[156:159], s[0:1] offset:624
+; CHECK-NEXT: global_store_dwordx4 v242, v[152:155], s[0:1] offset:608
+; CHECK-NEXT: global_store_dwordx4 v242, v[148:151], s[0:1] offset:592
+; CHECK-NEXT: global_store_dwordx4 v242, v[144:147], s[0:1] offset:576
+; CHECK-NEXT: global_store_dwordx4 v242, v[140:143], s[0:1] offset:560
+; CHECK-NEXT: global_store_dwordx4 v242, v[136:139], s[0:1] offset:544
+; CHECK-NEXT: global_store_dwordx4 v242, v[132:135], s[0:1] offset:528
+; CHECK-NEXT: global_store_dwordx4 v242, v[128:131], s[0:1] offset:512
+; CHECK-NEXT: global_store_dwordx4 v242, v[124:127], s[0:1] offset:752
+; CHECK-NEXT: global_store_dwordx4 v242, v[120:123], s[0:1] offset:736
+; CHECK-NEXT: global_store_dwordx4 v242, v[116:119], s[0:1] offset:720
+; CHECK-NEXT: global_store_dwordx4 v242, v[112:115], s[0:1] offset:704
+; CHECK-NEXT: global_store_dwordx4 v242, v[108:111], s[0:1] offset:688
+; CHECK-NEXT: global_store_dwordx4 v242, v[104:107], s[0:1] offset:672
+; CHECK-NEXT: global_store_dwordx4 v242, v[100:103], s[0:1] offset:656
+; CHECK-NEXT: global_store_dwordx4 v242, v[96:99], s[0:1] offset:640
+; CHECK-NEXT: global_store_dwordx4 v242, v[92:95], s[0:1] offset:880
+; CHECK-NEXT: global_store_dwordx4 v242, v[88:91], s[0:1] offset:864
+; CHECK-NEXT: global_store_dwordx4 v242, v[84:87], s[0:1] offset:848
+; CHECK-NEXT: global_store_dwordx4 v242, v[80:83], s[0:1] offset:832
+; CHECK-NEXT: global_store_dwordx4 v242, v[76:79], s[0:1] offset:816
+; CHECK-NEXT: global_store_dwordx4 v242, v[72:75], s[0:1] offset:800
+; CHECK-NEXT: global_store_dwordx4 v242, v[68:71], s[0:1] offset:784
+; CHECK-NEXT: global_store_dwordx4 v242, v[64:67], s[0:1] offset:768
+; CHECK-NEXT: global_store_dwordx4 v242, v[60:63], s[0:1] offset:1008
+; CHECK-NEXT: global_store_dwordx4 v242, v[56:59], s[0:1] offset:992
+; CHECK-NEXT: global_store_dwordx4 v242, v[52:55], s[0:1] offset:976
+; CHECK-NEXT: global_store_dwordx4 v242, v[48:51], s[0:1] offset:960
+; CHECK-NEXT: global_store_dwordx4 v242, v[44:47], s[0:1] offset:944
+; CHECK-NEXT: global_store_dwordx4 v242, v[40:43], s[0:1] offset:928
+; CHECK-NEXT: global_store_dwordx4 v242, v[36:39], s[0:1] offset:912
+; CHECK-NEXT: global_store_dwordx4 v242, v[32:35], s[0:1] offset:896
+; CHECK-NEXT: global_store_dwordx4 v242, v[28:31], s[0:1] offset:1136
+; CHECK-NEXT: global_store_dwordx4 v242, v[24:27], s[0:1] offset:1120
+; CHECK-NEXT: global_store_dwordx4 v242, v[20:23], s[0:1] offset:1104
+; CHECK-NEXT: global_store_dwordx4 v242, v[16:19], s[0:1] offset:1088
+; CHECK-NEXT: global_store_dwordx4 v242, v[12:15], s[0:1] offset:1072
+; CHECK-NEXT: global_store_dwordx4 v242, v[8:11], s[0:1] offset:1056
+; CHECK-NEXT: global_store_dwordx4 v242, v[4:7], s[0:1] offset:1040
+; CHECK-NEXT: global_store_dwordx4 v242, a[16:19], s[0:1] offset:1024
; CHECK-NEXT: s_endpgm
; CHECK-NEXT: .Ltmp2:
; Loop body: MFMAs in AGPR form.
diff --git a/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_copies.mir b/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_copies.mir
index a4a94089d2f7b..2f4cfe63693e4 100644
--- a/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_copies.mir
+++ b/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_copies.mir
@@ -2470,6 +2470,7 @@ body: |
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = COPY [[DS_READ_B128_gfx9_]]
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: $scc = IMPLICIT_DEF
@@ -2481,27 +2482,34 @@ body: |
; CHECK-NEXT: [[DEF14:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF15:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_1:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_2:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_4:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_5:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_4:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_5:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: KILL [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_1]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_2]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_4]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_5]]
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_5]]
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]]
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]]
+ ; CHECK-NEXT: [[COPY4:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_4]]
+ ; CHECK-NEXT: [[COPY5:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]]
+ ; CHECK-NEXT: [[COPY6:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]]
+ ; CHECK-NEXT: KILL [[COPY2]], [[COPY5]], [[COPY3]], [[COPY6]], [[COPY4]], [[COPY1]]
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.3:
; CHECK-NEXT: DS_WRITE_B128_gfx9 [[DEF11]], [[DS_READ_B128_gfx9_]], 0, 0, implicit $exec
; CHECK-NEXT: SCHED_BARRIER 0
+ ; CHECK-NEXT: [[COPY7:%[0-9]+]]:vreg_128_align2 = COPY [[COPY]]
; CHECK-NEXT: [[DEF17:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
- ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF12]], [[DEF13]], [[DS_READ_B128_gfx9_]]
+ ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF12]], [[DEF13]], [[COPY7]]
; CHECK-NEXT: S_NOP 0, implicit %12, implicit %13
; CHECK-NEXT: S_ENDPGM 0
bb.0:
@@ -2592,6 +2600,7 @@ body: |
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = COPY [[DS_READ_B128_gfx9_]]
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: $scc = IMPLICIT_DEF
@@ -2603,27 +2612,34 @@ body: |
; CHECK-NEXT: [[DEF14:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF15:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_1:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_2:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_4:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_5:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_4:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_5:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: KILL [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_1]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_2]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_4]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_5]]
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_5]]
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]]
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]]
+ ; CHECK-NEXT: [[COPY4:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_4]]
+ ; CHECK-NEXT: [[COPY5:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]]
+ ; CHECK-NEXT: [[COPY6:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]]
+ ; CHECK-NEXT: KILL [[COPY2]], [[COPY5]], [[COPY3]], [[COPY6]], [[COPY4]], [[COPY1]]
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.3:
- ; CHECK-NEXT: [[V_ADD_U32_e32_:%[0-9]+]]:vgpr_32 = V_ADD_U32_e32 [[DEF13]].sub1, [[DS_READ_B128_gfx9_]].sub0, implicit $exec
+ ; CHECK-NEXT: [[COPY7:%[0-9]+]]:vreg_128_align2 = COPY [[COPY]]
+ ; CHECK-NEXT: [[V_ADD_U32_e32_:%[0-9]+]]:vgpr_32 = V_ADD_U32_e32 [[DEF13]].sub1, [[COPY7]].sub0, implicit $exec
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF17:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
- ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF12]], [[DEF13]], [[DS_READ_B128_gfx9_]], [[V_ADD_U32_e32_]]
+ ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF12]], [[DEF13]], [[COPY7]], [[V_ADD_U32_e32_]]
; CHECK-NEXT: S_NOP 0, implicit %12, implicit %13
; CHECK-NEXT: S_ENDPGM 0
bb.0:
@@ -3560,7 +3576,7 @@ body: |
; CHECK-NEXT: [[DEF10:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: {{ $}}
@@ -3570,10 +3586,10 @@ body: |
; CHECK-NEXT: [[DEF14:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF15:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_1:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_2:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_1]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_2]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.4(0x40000000), %bb.3(0x40000000)
@@ -3584,21 +3600,25 @@ body: |
; CHECK-NEXT: bb.3:
; CHECK-NEXT: successors: %bb.5(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: DS_WRITE_B32_gfx9 [[DEF11]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3]].sub0, 0, 0, implicit $exec
- ; CHECK-NEXT: DS_WRITE_B32_gfx9 [[DEF11]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3]].sub1, 256, 0, implicit $exec
+ ; CHECK-NEXT: DS_WRITE_B32_gfx9 [[DEF11]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]].sub0, 0, 0, implicit $exec
+ ; CHECK-NEXT: DS_WRITE_B32_gfx9 [[DEF11]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]].sub1, 256, 0, implicit $exec
; CHECK-NEXT: S_BRANCH %bb.5
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.4:
; CHECK-NEXT: successors: %bb.5(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: DS_WRITE_B32_gfx9 [[DEF11]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3]].sub1, 0, 0, implicit $exec
- ; CHECK-NEXT: DS_WRITE_B32_gfx9 [[DEF11]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3]].sub0, 256, 0, implicit $exec
+ ; CHECK-NEXT: DS_WRITE_B32_gfx9 [[DEF11]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]].sub1, 0, 0, implicit $exec
+ ; CHECK-NEXT: DS_WRITE_B32_gfx9 [[DEF11]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]].sub0, 256, 0, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.5:
; CHECK-NEXT: DS_WRITE_B128_gfx9 [[DEF11]], [[DS_READ_B128_gfx9_]], 0, 0, implicit $exec
; CHECK-NEXT: SCHED_BARRIER 0
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]]
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]]
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]]
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]]
; CHECK-NEXT: [[DEF17:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
- ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF12]], [[DEF13]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_1]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_2]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3]]
+ ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF12]], [[DEF13]], [[COPY]], [[COPY2]], [[COPY1]], [[COPY3]]
; CHECK-NEXT: S_NOP 0, implicit %12, implicit %13
; CHECK-NEXT: S_ENDPGM 0
bb.0:
@@ -5479,18 +5499,16 @@ body: |
; CHECK-NEXT: [[DEF15:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DEF13]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_]]
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_1:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_2:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_1]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_2]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: undef [[V_ADD_U32_e32_:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_]].sub0, [[DEF12]], implicit $exec
- ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]]
- ; CHECK-NEXT: [[COPY2:%[0-9]+]]:vreg_128_align2 = COPY [[COPY1]]
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3]]
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF17:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
- ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF11]], [[DEF13]], [[V_ADD_U32_e32_]], [[COPY2]]
+ ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF11]], [[DEF13]], [[V_ADD_U32_e32_]], [[COPY]]
; CHECK-NEXT: S_NOP 0, implicit %12, implicit %13
; CHECK-NEXT: S_ENDPGM 0
bb.0:
@@ -5580,10 +5598,9 @@ body: |
; CHECK-NEXT: [[DEF10:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
; CHECK-NEXT: undef [[V_ADD_U32_e32_:%[0-9]+]].sub0:vreg_64_align2 = V_ADD_U32_e32 [[AV_MOV_]].sub0, [[DEF10]], implicit $exec
; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_64_align2 = COPY [[V_MFMA_F64_4X4X4F64_e64_3]]
- ; CHECK-NEXT: [[COPY2:%[0-9]+]]:vreg_64_align2 = COPY [[COPY1]]
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
- ; CHECK-NEXT: KILL [[DEF11]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[V_ADD_U32_e32_]], [[COPY2]]
+ ; CHECK-NEXT: KILL [[DEF11]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[V_ADD_U32_e32_]], [[COPY1]]
; CHECK-NEXT: S_NOP 0, implicit %9, implicit %10
; CHECK-NEXT: S_ENDPGM 0
bb.0:
>From 1a908ac02a965e8707c276244cc7a7efa5c19864 Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Fri, 21 Aug 2026 13:54:02 +0000
Subject: [PATCH 2/9] [AMDGPU] NFC: simplify isRecolorSafe reaching-def scan
Walk def_instructions instead of def_operands so a def is visited once
per instruction (dropping the getParent hop and avoiding duplicate
findReachingUses work when one instr partially defs multiple subregs),
and fold the two remaining raw isMAI() && RewriteSet.contains() checks
into the existing isRewriteCandidateMAI helper.
Co-Authored-By: Claude Opus 4.7 <noreply at anthropic.com>
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 13 ++++++-------
1 file changed, 6 insertions(+), 7 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index 59d4ae7ebadb2..542861d552017 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -2350,17 +2350,16 @@ static bool useAcceptsAGPR(const MachineOperand *Use, const SIInstrInfo *TII,
bool RewriteMFMAFormStage::isRecolorSafe(
Register Reg, ArrayRef<MachineOperand *> DstReachingUses,
const SmallPtrSetImpl<MachineInstr *> &RewriteCandsSet, bool IsDst) {
- for (MachineOperand &DefMO : DAG.MRI.def_operands(Reg)) {
- MachineInstr *DefMI = DefMO.getParent();
+ for (MachineInstr &DefMI : DAG.MRI.def_instructions(Reg)) {
// A candidate MFMA def is rewritten to AGPR form (it produces the AGPR
// result directly), so it does not constrain the recolor.
- if (isRewriteCandidateMAI(DefMI, TII, RewriteCandsSet))
+ if (isRewriteCandidateMAI(&DefMI, TII, RewriteCandsSet))
continue;
- bool CanWriteAGPR = canWriteAGPR(DefMI, TII, SRI);
+ bool CanWriteAGPR = canWriteAGPR(&DefMI, TII, SRI);
if (!CanWriteAGPR)
return false;
SmallVector<MachineOperand *, 8> DefReachingUses;
- findReachingUses(DefMI, DAG.LIS, DefReachingUses);
+ findReachingUses(&DefMI, DAG.LIS, DefReachingUses);
for (MachineOperand *UseMO : DefReachingUses) {
// Exempt uses that do not constrain the recolor: for a dst, its own
// reaching uses (bridged to VGPR by case2); for a src2, uses that are
@@ -2466,7 +2465,7 @@ bool RewriteMFMAFormStage::initHeuristics(
MachineInstr *UserMI = RUOp->getParent();
bool NeedsAGPRToVGPRCopy = true;
// Group members read the AGPR result directly.
- if (TII->isMAI(*UserMI) && RewriteSet.contains(UserMI))
+ if (isRewriteCandidateMAI(UserMI, TII, RewriteSet))
NeedsAGPRToVGPRCopy = false;
// For any user of the result of the MFMA which is not an MFMA, we
@@ -2484,7 +2483,7 @@ bool RewriteMFMAFormStage::initHeuristics(
for (SlotIndex RDIndex : DstUsesReachingDefs) {
MachineInstr *RD = DAG.LIS->getInstructionFromIndex(RDIndex);
- if (TII->isMAI(*RD) && RewriteSet.contains(RD))
+ if (isRewriteCandidateMAI(RD, TII, RewriteSet))
continue;
// Dst cannot be recolored: for any non-MFMA reaching def of this
>From cadce1f2b1e12cdcf48706f3cf128ad1d9cf306b Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Tue, 25 Aug 2026 10:44:21 +0000
Subject: [PATCH 3/9] [AMDGPU] Factor operand AGPR-form check into
operandAcceptsAGPRForm
Factor the duplicated "does this operand's RC admit the AGPR form we
emit" test out of canWriteAGPR and useAcceptsAGPR into a shared
operandAcceptsAGPRForm helper.
While here, fix a verifier error in RewriteMFMAFormStage::rewrite(): an
unsafe, non-recolorable MFMA dst with no reaching use was left without a
full-lane def after its MFMA def moved to a mapped AGPR reg. Bridge the
AGPR-form result back into the original VGPR dst with a COPY.
Co-Authored-By: Claude Opus 4.7 <noreply at anthropic.com>
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 38 +++++--
...ewrite-mfma-form-unsafe-recolor-no-use.mir | 105 ++++++++++++++++++
2 files changed, 133 insertions(+), 10 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-recolor-no-use.mir
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index 542861d552017..6f731dc5f450c 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -2320,19 +2320,27 @@ isRewriteCandidateMAI(const MachineInstr *MI, const SIInstrInfo *TII,
return TII->isMAI(*MI) && RewriteCandsSet.contains(MI);
}
+// True if \p OpRC (an operand's register-class constraint) admits the AGPR
+// form we would emit for it, i.e. the AGPR equivalent of its VGPR class is a
+// legal choice. An AV/agnostic constraint passes; a plain-VGPR one does not.
+static bool operandAcceptsAGPRForm(const TargetRegisterClass *OpRC,
+ const SIRegisterInfo *SRI) {
+ if (!OpRC)
+ return false;
+ const TargetRegisterClass *AGPRForm = SRI->getEquivalentAGPRClass(OpRC);
+ return AGPRForm && SRI->getCommonSubClass(OpRC, AGPRForm);
+}
+
static bool canWriteAGPR(const MachineInstr *MI, const SIInstrInfo *TII,
const SIRegisterInfo *SRI) {
- // A writer can write its result into an AGPR lane iff its def operand's
- // register-class constraint admits AGPRs (AV/agnostic classes do, plain VGPR
- // classes don't). COPY is a generic opcode with no operand constraint, so
- // special-case it.
+ // COPY is a generic opcode with no operand constraint, so it can produce an
+ // AGPR result.
if (MI->isCopy())
return true;
const MCInstrDesc &Desc = MI->getDesc();
if (Desc.getNumDefs() == 0)
return false;
- const TargetRegisterClass *DefRC = TII->getRegClass(Desc, 0);
- return DefRC && SRI->hasAGPRs(DefRC);
+ return operandAcceptsAGPRForm(TII->getRegClass(Desc, 0), SRI);
}
static bool useAcceptsAGPR(const MachineOperand *Use, const SIInstrInfo *TII,
@@ -2341,10 +2349,8 @@ static bool useAcceptsAGPR(const MachineOperand *Use, const SIInstrInfo *TII,
// COPY is a generic opcode with no operand constraint; it can read an AGPR.
if (UseMI->isCopy())
return true;
- // The use's static register-class constraint must admit AGPRs.
- const TargetRegisterClass *UseRC =
- UseMI->getRegClassConstraint(Use->getOperandNo(), TII, SRI);
- return UseRC && SRI->hasAGPRs(UseRC);
+ return operandAcceptsAGPRForm(
+ UseMI->getRegClassConstraint(Use->getOperandNo(), TII, SRI), SRI);
}
bool RewriteMFMAFormStage::isRecolorSafe(
@@ -2830,6 +2836,18 @@ bool RewriteMFMAFormStage::rewrite(
const TargetRegisterClass *VGPRRC = SRI->getEquivalentVGPRClass(DstRC);
MappedReg = DAG.MRI.createVirtualRegister(VGPRRC);
RedefMap[DstReg] = MappedReg;
+
+ // The candidate MFMA's def operand is redirected to MappedReg (which is
+ // reclassified to AGPR form), so the original dst reg loses its full-lane
+ // def. Bridge the AGPR-form result back into the original VGPR reg right
+ // after the MFMA so any later partial redef and downstream use still see
+ // a dominating full-lane def.
+ MachineInstrBuilder Bridge =
+ BuildMI(*MI->getParent(), std::next(MI->getIterator()),
+ MI->getDebugLoc(), TII->get(TargetOpcode::COPY))
+ .addDef(DstReg, {}, 0)
+ .addUse(MappedReg, {}, 0);
+ DAG.LIS->InsertMachineInstrInMaps(*Bridge);
}
if (!DstUseDefsReplace.empty()) {
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-recolor-no-use.mir b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-recolor-no-use.mir
new file mode 100644
index 0000000000000..b3deb844fb28e
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-recolor-no-use.mir
@@ -0,0 +1,105 @@
+# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py
+# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -run-pass=machine-scheduler -amdgpu-disable-rewrite-mfma-form-sched-stage=false -verify-machineinstrs -o - %s | FileCheck %s
+#
+# Covers the "DstReachingUses.empty() && !DstRecolorSafe" branch in
+# RewriteMFMAFormStage::rewrite().
+
+--- |
+ define void @cover2831() #0 { entry: unreachable }
+ attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }
+...
+---
+name: cover2831
+tracksRegLiveness: true
+machineFunctionInfo:
+ isEntryFunction: true
+ scratchRSrcReg: '$sgpr96_sgpr97_sgpr98_sgpr99'
+ stackPtrOffsetReg: '$sgpr32'
+ argumentInfo:
+ privateSegmentBuffer: { reg: '$sgpr0_sgpr1_sgpr2_sgpr3' }
+ kernargSegmentPtr: { reg: '$sgpr4_sgpr5' }
+ workGroupIDX: { reg: '$sgpr6' }
+ privateSegmentWaveByteOffset: { reg: '$sgpr7' }
+ workItemIDX: { reg: '$vgpr0' }
+ sgprForEXECCopy: '$sgpr100_sgpr101'
+body: |
+ ; CHECK-LABEL: name: cover2831
+ ; CHECK: bb.0:
+ ; CHECK-NEXT: successors: %bb.1(0x80000000)
+ ; CHECK-NEXT: liveins: $vgpr0, $sgpr4_sgpr5
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[DEF:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF1:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF2:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF3:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF4:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF5:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF6:%[0-9]+]]:vreg_512 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF7:%[0-9]+]]:vreg_64 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF8:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF9:%[0-9]+]]:vreg_128 = IMPLICIT_DEF
+ ; CHECK-NEXT: S_NOP 0, implicit-def %12
+ ; CHECK-NEXT: S_NOP 0, implicit-def %13
+ ; CHECK-NEXT: [[DEF10:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: SCHED_BARRIER 0
+ ; CHECK-NEXT: [[DEF11:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF12:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = COPY [[DEF12]]
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: bb.1:
+ ; CHECK-NEXT: successors: %bb.2(0x80000000)
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[DEF13:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF13]], [[DEF14]], [[COPY]], 4, 4, [[DEF15]].sub0, [[DEF16]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF13]], [[DEF14]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]], 4, 4, [[DEF15]].sub0, [[DEF16]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF13]], [[DEF14]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]], 4, 4, [[DEF15]].sub0, [[DEF16]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF13]], [[DEF14]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]], 4, 4, [[DEF15]].sub0, [[DEF16]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]]
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[DEF11]].sub0, [[DEF16]], implicit $exec
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: bb.2:
+ ; CHECK-NEXT: SCHED_BARRIER 0
+ ; CHECK-NEXT: [[DEF17:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF11]], [[DEF12]], [[COPY1]]
+ ; CHECK-NEXT: S_NOP 0, implicit %12, implicit %13
+ ; CHECK-NEXT: S_ENDPGM 0
+ bb.0:
+ liveins: $vgpr0, $sgpr4_sgpr5
+ %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_512 = IMPLICIT_DEF
+ %8:vreg_64 = IMPLICIT_DEF
+ %9:vgpr_32 = IMPLICIT_DEF
+ %10:vreg_128 = IMPLICIT_DEF
+ %11:vreg_1024 = IMPLICIT_DEF
+ S_NOP 0, implicit-def %12:av_512
+ S_NOP 0, implicit-def %13:av_512
+ SCHED_BARRIER 0
+ %14:av_128_align2 = IMPLICIT_DEF
+ %15:av_128_align2 = IMPLICIT_DEF
+ %16:vreg_128_align2 = IMPLICIT_DEF
+ %17:vreg_64_align2 = IMPLICIT_DEF
+ %18:vgpr_32 = IMPLICIT_DEF
+ %19:vreg_128_align2 = IMPLICIT_DEF
+
+ bb.1:
+ %20:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 %14, %15, %19, 4, 4, %17.sub0, %18, 0, 0, implicit $mode, implicit $exec
+ %21:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 %14, %15, %20, 4, 4, %17.sub0, %18, 0, 0, implicit $mode, implicit $exec
+ %22:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 %14, %15, %21, 4, 4, %17.sub0, %18, 0, 0, implicit $mode, implicit $exec
+ %23:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 %14, %15, %22, 4, 4, %17.sub0, %18, 0, 0, implicit $mode, implicit $exec
+ %23.sub0:vreg_128_align2 = V_ADD_U32_e32 %16.sub0, %18, implicit $exec
+
+ bb.2:
+ SCHED_BARRIER 0
+ KILL %0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11, %16, %19, %23
+ S_NOP 0, implicit %12, implicit %13
+ S_ENDPGM 0
+...
>From a9fe4a23dc8d66e4468ddbdd915d3e51e87e56d3 Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Thu, 27 Aug 2026 06:02:58 +0000
Subject: [PATCH 4/9] [AMDGPU] Map recolor-unsafe MFMA dst even without a
bridge def
A recolor-unsafe candidate MFMA dst that has reaching uses but no
non-candidate reaching def to bridge (DstUseDefsReplace empty) was
never recorded in the redef map, so the final reclassification recolored
the original VGPR dst to AGPR instead of a fresh carrier.
Record the mapping for this case so the original dst stays VGPR and only
the mapped AGPR carrier is reclassified. Adds a dedicated test and
updates sched_mfma_rewrite_copies.mir to store the correct value.
Co-Authored-By: Claude Opus 4.7 <noreply at anthropic.com>
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 38 +++++-
...form-unsafe-dst-no-bridge-still-mapped.mir | 124 ++++++++++++++++++
.../AMDGPU/sched_mfma_rewrite_copies.mir | 4 +-
3 files changed, 160 insertions(+), 6 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-dst-no-bridge-still-mapped.mir
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index 6f731dc5f450c..531a907ae88cd 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -2804,10 +2804,18 @@ bool RewriteMFMAFormStage::rewrite(
isRecolorSafe(DstReg, DstReachingUses, RewriteCandsSet, /*IsDst=*/true);
for (MachineOperand *RUOp : DstReachingUses) {
MachineInstr *UserMI = RUOp->getParent();
- // A group-member MFMA reads the AGPR result directly; any other user can
- // only avoid a bridge if it accepts an AGPR operand.
- bool CanReadAGPR = TII->isMAI(*UserMI) ? RewriteCandsSet.contains(UserMI)
- : useAcceptsAGPR(RUOp, TII, SRI);
+ // Decide whether this reaching use can read the dst's AGPR form directly
+ // or needs an AGPR->VGPR bridge copy.
+ // - A group-member MFMA always reads the AGPR result directly.
+ // - Any other user can skip the bridge only when the dst is recolored
+ // to AGPR (DstRecolorSafe) and its operand accepts an AGPR. When the
+ // dst is unsafe, its original reg stays VGPR, so every non-MFMA user
+ // must go through a bridge copy.
+ bool CanReadAGPR;
+ if (TII->isMAI(*UserMI))
+ CanReadAGPR = RewriteCandsSet.contains(UserMI);
+ else
+ CanReadAGPR = DstRecolorSafe && useAcceptsAGPR(RUOp, TII, SRI);
if (!CanReadAGPR &&
find(DstReachingUseCopies, RUOp) == DstReachingUseCopies.end())
DstReachingUseCopies.push_back(RUOp);
@@ -2832,6 +2840,11 @@ bool RewriteMFMAFormStage::rewrite(
// reg to carry the AGPR-form value and record the mapping, leaving the
// original dst reg in VGPR form.
if (DstReachingUses.empty() && !DstRecolorSafe) {
+ // Exclusion must already have dropped any dst that was bridged as an
+ // earlier MFMA's src2, so it cannot be pre-mapped when we reach here.
+ assert(
+ !RedefMap.contains(DstReg) &&
+ "empty-use dst unexpectedly already mapped -- exclusion missed it");
const TargetRegisterClass *DstRC = DAG.MRI.getRegClass(DstReg);
const TargetRegisterClass *VGPRRC = SRI->getEquivalentVGPRClass(DstRC);
MappedReg = DAG.MRI.createVirtualRegister(VGPRRC);
@@ -2850,6 +2863,23 @@ bool RewriteMFMAFormStage::rewrite(
DAG.LIS->InsertMachineInstrInMaps(*Bridge);
}
+ // Unsafe dst with no def to bridge: no copy needed, but it must still be
+ // mapped so reclassification recolors the mapped reg instead of illegally
+ // recoloring the original VGPR dst to AGPR.
+ if (DstUseDefsReplace.empty() && !DstRecolorSafe) {
+ auto RI = RedefMap.find(DstReg);
+ if (RI != RedefMap.end()) {
+ MappedReg = RI->second;
+ } else {
+ assert(!ReachingDefCopyMap.contains(DstReg));
+ const TargetRegisterClass *DstRC = DAG.MRI.getRegClass(DstReg);
+ const TargetRegisterClass *VGPRRC = SRI->getEquivalentVGPRClass(DstRC);
+
+ // Track the mapping of the original register to the new register.
+ MappedReg = DAG.MRI.createVirtualRegister(VGPRRC);
+ RedefMap[DstReg] = MappedReg;
+ }
+ }
if (!DstUseDefsReplace.empty()) {
auto RI = RedefMap.find(DstReg);
if (RI != RedefMap.end()) {
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-dst-no-bridge-still-mapped.mir b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-dst-no-bridge-still-mapped.mir
new file mode 100644
index 0000000000000..2aeaffc8f537d
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-dst-no-bridge-still-mapped.mir
@@ -0,0 +1,124 @@
+# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 6
+# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a \
+# RUN: -run-pass=machine-scheduler \
+# RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+# RUN: -verify-machineinstrs %s -o - | FileCheck %s
+#
+# A candidate MFMA whose dst %acc is recolor-unsafe but which has NO
+# non-candidate reaching-use def (DstUseDefsReplace is empty), must still be
+# recorded in the redef map so the dst reg stays VGPR and only a fresh AGPR
+# carrier is reclassified. Without that mapping the dst reg is left in the
+# reclassification
+# set and gets wholly recolored to AGPR -- illegal, because the whole-register
+# V_ADD_F64_e64 redef is a VALU write that cannot target an AReg subregister
+# ("Expected a VReg_64_Align2 register, but got a AReg_64_Align2 register").
+#
+# CFG:
+# bb.0 %acc = MFMA(...) (candidate dst)
+# / \
+# bb.1 bb.3 -> bb.2 (self-loop) bb.2: MFMA reads %acc as src2
+# | \ (loop inflates spill savings so the
+# %acc = V_ADD %acc bb.4 cost gate is beneficial)
+# (unsafe whole-reg redef, unreachable to the loop use)
+--- |
+ define void @unsafe_dst_no_bridge_still_mapped() #0 { ret void }
+ attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }
+...
+---
+name: unsafe_dst_no_bridge_still_mapped
+tracksRegLiveness: true
+body: |
+ ; CHECK-LABEL: name: unsafe_dst_no_bridge_still_mapped
+ ; CHECK: bb.0:
+ ; CHECK-NEXT: successors: %bb.1(0x40000000), %bb.2(0x40000000)
+ ; CHECK-NEXT: liveins: $vgpr0, $sgpr4_sgpr5, $sgpr6
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: %cnt:sgpr_32 = COPY $sgpr6
+ ; CHECK-NEXT: %s0:vreg_64_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: %s1:vreg_64_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_e64_:%[0-9]+]]:areg_64_align2 = nofpexcept V_MFMA_F64_4X4X4F64_e64 %s0, %s1, 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %p0:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p1:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p2:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p3:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p4:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p5:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p6:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p7:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p8:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p9:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p10:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p11:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p12:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: SCHED_BARRIER 0
+ ; CHECK-NEXT: S_CBRANCH_SCC1 %bb.1, implicit undef $scc
+ ; CHECK-NEXT: S_BRANCH %bb.2
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: bb.1:
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY [[V_MFMA_F64_4X4X4F64_e64_]]
+ ; CHECK-NEXT: dead %acc:vreg_64_align2 = nofpexcept V_ADD_F64_e64 0, [[COPY]], 0, %s0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: S_ENDPGM 0, implicit %p0, implicit %p1, implicit %p2, implicit %p3, implicit %p4, implicit %p5, implicit %p6, implicit %p7, implicit %p8, implicit %p9, implicit %p10, implicit %p11, implicit %p12
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: bb.2:
+ ; CHECK-NEXT: successors: %bb.3(0x80000000)
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: S_BRANCH %bb.3
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: bb.3:
+ ; CHECK-NEXT: successors: %bb.3(0x40000000), %bb.4(0x40000000)
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: dead %mdst:areg_64_align2 = nofpexcept V_MFMA_F64_4X4X4F64_e64 %s0, %s1, [[V_MFMA_F64_4X4X4F64_e64_]], 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %cnt:sgpr_32 = S_SUB_I32 %cnt, 1, implicit-def $scc
+ ; CHECK-NEXT: S_CBRANCH_SCC1 %bb.3, implicit $scc
+ ; CHECK-NEXT: S_BRANCH %bb.4
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: bb.4:
+ ; CHECK-NEXT: KILL %p0, %p1, %p2, %p3, %p4, %p5, %p6, %p7, %p8, %p9, %p10, %p11, %p12
+ ; CHECK-NEXT: S_ENDPGM 0
+ bb.0:
+ successors: %bb.1(0x40000000), %bb.3(0x40000000)
+ liveins: $vgpr0, $sgpr4_sgpr5, $sgpr6
+
+ %p0:vreg_1024 = IMPLICIT_DEF
+ %p1:vreg_1024 = IMPLICIT_DEF
+ %p2:vreg_1024 = IMPLICIT_DEF
+ %p3:vreg_1024 = IMPLICIT_DEF
+ %p4:vreg_1024 = IMPLICIT_DEF
+ %p5:vreg_1024 = IMPLICIT_DEF
+ %p6:vreg_1024 = IMPLICIT_DEF
+ %p7:vreg_1024 = IMPLICIT_DEF
+ %p8:vreg_1024 = IMPLICIT_DEF
+ %p9:vreg_1024 = IMPLICIT_DEF
+ %p10:vreg_1024 = IMPLICIT_DEF
+ %p11:vreg_1024 = IMPLICIT_DEF
+ %p12:vreg_1024 = IMPLICIT_DEF
+ %s0:vreg_64_align2 = IMPLICIT_DEF
+ %s1:vreg_64_align2 = IMPLICIT_DEF
+ %cnt:sgpr_32 = COPY $sgpr6
+ %acc:vreg_64_align2 = nofpexcept V_MFMA_F64_4X4X4F64_vgprcd_e64 %s0, %s1, 0, 0, 0, 0, implicit $mode, implicit $exec
+ SCHED_BARRIER 0
+ S_CBRANCH_SCC1 %bb.1, implicit undef $scc
+ S_BRANCH %bb.3
+
+ bb.1:
+ ; Unsafe whole-register def of %acc (non-candidate VALU, cannot write AGPR).
+ %acc:vreg_64_align2 = nofpexcept V_ADD_F64_e64 0, %acc, 0, %s0, 0, 0, implicit $mode, implicit $exec
+ S_ENDPGM 0, implicit %p0, implicit %p1, implicit %p2, implicit %p3, implicit %p4, implicit %p5, implicit %p6, implicit %p7, implicit %p8, implicit %p9, implicit %p10, implicit %p11, implicit %p12
+
+ bb.3:
+ successors: %bb.2
+ S_BRANCH %bb.2
+
+ bb.2:
+ ; Self-loop MFMA reading %acc as src2; the loop inflates the spill savings so
+ ; the profitability gate accepts the rewrite.
+ successors: %bb.2(0x40000000), %bb.4(0x40000000)
+ dead %mdst:vreg_64_align2 = nofpexcept V_MFMA_F64_4X4X4F64_vgprcd_e64 %s0, %s1, %acc, 0, 0, 0, implicit $mode, implicit $exec
+ %cnt:sgpr_32 = S_SUB_I32 %cnt, 1, implicit-def $scc
+ S_CBRANCH_SCC1 %bb.2, implicit $scc
+ S_BRANCH %bb.4
+
+ bb.4:
+ KILL %p0, %p1, %p2, %p3, %p4, %p5, %p6, %p7, %p8, %p9, %p10, %p11, %p12
+ S_ENDPGM 0
+...
diff --git a/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_copies.mir b/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_copies.mir
index 2f4cfe63693e4..845ab48af9ab2 100644
--- a/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_copies.mir
+++ b/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_copies.mir
@@ -2505,9 +2505,9 @@ body: |
; CHECK-NEXT: KILL [[COPY2]], [[COPY5]], [[COPY3]], [[COPY6]], [[COPY4]], [[COPY1]]
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.3:
- ; CHECK-NEXT: DS_WRITE_B128_gfx9 [[DEF11]], [[DS_READ_B128_gfx9_]], 0, 0, implicit $exec
- ; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[COPY7:%[0-9]+]]:vreg_128_align2 = COPY [[COPY]]
+ ; CHECK-NEXT: DS_WRITE_B128_gfx9 [[DEF11]], [[COPY7]], 0, 0, implicit $exec
+ ; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF17:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF12]], [[DEF13]], [[COPY7]]
; CHECK-NEXT: S_NOP 0, implicit %12, implicit %13
>From a672d7a08b4416c509355badc95f8e1d95520160 Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Thu, 27 Aug 2026 12:34:57 +0000
Subject: [PATCH 5/9] [AMDGPU] Use folded gfx triple in rewrite-mfma-form
recolor tests
Switch the RUN lines to the -mtriple=amdgpuX.Y-amd-amdhsa form (mcpu
folded into the triple) to match the sibling rewrite-mfma-form tests.
Co-Authored-By: Claude Opus 4.7 <noreply at anthropic.com>
---
.../rewrite-mfma-form-unsafe-dst-no-bridge-still-mapped.mir | 2 +-
.../CodeGen/AMDGPU/rewrite-mfma-form-unsafe-recolor-no-use.mir | 2 +-
2 files changed, 2 insertions(+), 2 deletions(-)
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-dst-no-bridge-still-mapped.mir b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-dst-no-bridge-still-mapped.mir
index 2aeaffc8f537d..79e98e75fc7ac 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-dst-no-bridge-still-mapped.mir
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-dst-no-bridge-still-mapped.mir
@@ -1,5 +1,5 @@
# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 6
-# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a \
+# RUN: llc -mtriple=amdgpu9.0a-amd-amdhsa \
# RUN: -run-pass=machine-scheduler \
# RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
# RUN: -verify-machineinstrs %s -o - | FileCheck %s
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-recolor-no-use.mir b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-recolor-no-use.mir
index b3deb844fb28e..83cc85686fb15 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-recolor-no-use.mir
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-recolor-no-use.mir
@@ -1,5 +1,5 @@
# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py
-# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -run-pass=machine-scheduler -amdgpu-disable-rewrite-mfma-form-sched-stage=false -verify-machineinstrs -o - %s | FileCheck %s
+# RUN: llc -mtriple=amdgpu9.50-amd-amdhsa -run-pass=machine-scheduler -amdgpu-disable-rewrite-mfma-form-sched-stage=false -verify-machineinstrs -o - %s | FileCheck %s
#
# Covers the "DstReachingUses.empty() && !DstRecolorSafe" branch in
# RewriteMFMAFormStage::rewrite().
>From b655979dc04e90cead8ae0b377ba524e5b47ee9b Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Thu, 3 Sep 2026 03:00:21 +0000
Subject: [PATCH 6/9] [AMDGPU] Use getRegClassConstraintEffect for MFMA recolor
legality
Reform canWriteAGPR/useAcceptsAGPR to test AGPR legality via
MachineInstr::getRegClassConstraintEffect against the target AGPR class,
mirroring the def/use handling in AMDGPURewriteAGPRCopyMFMA. This folds
in each operand's subregister index and correctly treats unconstrained
defs (e.g. IMPLICIT_DEF) as AGPR-placeable, replacing the
over-conservative operandAcceptsAGPRForm helper (now removed). Regenerate
the affected flag-test CHECKs.
Co-Authored-By: Claude Opus 4.7 <noreply at anthropic.com>
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 55 ++--
.../AMDGPU/mfma-no-multiple-copies.mir | 9 +-
...ewrite-mfma-form-unsafe-recolor-no-use.mir | 11 +-
.../AMDGPU/sched_mfma_rewrite_copies.mir | 280 +++++++++---------
.../AMDGPU/sched_mfma_rewrite_cost.mir | 16 +-
.../AMDGPU/sched_mfma_rewrite_diff_types.mir | 6 +-
6 files changed, 180 insertions(+), 197 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index 531a907ae88cd..166500c02630c 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -2320,49 +2320,38 @@ isRewriteCandidateMAI(const MachineInstr *MI, const SIInstrInfo *TII,
return TII->isMAI(*MI) && RewriteCandsSet.contains(MI);
}
-// True if \p OpRC (an operand's register-class constraint) admits the AGPR
-// form we would emit for it, i.e. the AGPR equivalent of its VGPR class is a
-// legal choice. An AV/agnostic constraint passes; a plain-VGPR one does not.
-static bool operandAcceptsAGPRForm(const TargetRegisterClass *OpRC,
- const SIRegisterInfo *SRI) {
- if (!OpRC)
+static bool canWriteAGPR(const MachineInstr *MI, Register Reg,
+ const TargetRegisterClass *RegAGPRClass,
+ const SIInstrInfo *TII, const SIRegisterInfo *SRI) {
+ if (MI->getDesc().getNumDefs() == 0 || !RegAGPRClass)
return false;
- const TargetRegisterClass *AGPRForm = SRI->getEquivalentAGPRClass(OpRC);
- return AGPRForm && SRI->getCommonSubClass(OpRC, AGPRForm);
+ int DefOpIdx =
+ MI->findRegisterDefOperandIdx(Reg, /*TRI=*/nullptr, false, false);
+ return DefOpIdx >= 0 && MI->getRegClassConstraintEffect(
+ DefOpIdx, RegAGPRClass, TII, SRI) != nullptr;
}
-static bool canWriteAGPR(const MachineInstr *MI, const SIInstrInfo *TII,
- const SIRegisterInfo *SRI) {
- // COPY is a generic opcode with no operand constraint, so it can produce an
- // AGPR result.
- if (MI->isCopy())
- return true;
- const MCInstrDesc &Desc = MI->getDesc();
- if (Desc.getNumDefs() == 0)
- return false;
- return operandAcceptsAGPRForm(TII->getRegClass(Desc, 0), SRI);
-}
-
-static bool useAcceptsAGPR(const MachineOperand *Use, const SIInstrInfo *TII,
- const SIRegisterInfo *SRI) {
+static bool useAcceptsAGPR(const MachineOperand *Use,
+ const TargetRegisterClass *RegAGPRClass,
+ const SIInstrInfo *TII, const SIRegisterInfo *SRI) {
const MachineInstr *UseMI = Use->getParent();
- // COPY is a generic opcode with no operand constraint; it can read an AGPR.
- if (UseMI->isCopy())
- return true;
- return operandAcceptsAGPRForm(
- UseMI->getRegClassConstraint(Use->getOperandNo(), TII, SRI), SRI);
+ if (!RegAGPRClass)
+ return false;
+ return UseMI->getRegClassConstraintEffect(Use->getOperandNo(), RegAGPRClass,
+ TII, SRI) != nullptr;
}
bool RewriteMFMAFormStage::isRecolorSafe(
Register Reg, ArrayRef<MachineOperand *> DstReachingUses,
const SmallPtrSetImpl<MachineInstr *> &RewriteCandsSet, bool IsDst) {
+ const TargetRegisterClass *RegAGPRClass =
+ SRI->getEquivalentAGPRClass(DAG.MRI.getRegClass(Reg));
for (MachineInstr &DefMI : DAG.MRI.def_instructions(Reg)) {
// A candidate MFMA def is rewritten to AGPR form (it produces the AGPR
// result directly), so it does not constrain the recolor.
if (isRewriteCandidateMAI(&DefMI, TII, RewriteCandsSet))
continue;
- bool CanWriteAGPR = canWriteAGPR(&DefMI, TII, SRI);
- if (!CanWriteAGPR)
+ if (!canWriteAGPR(&DefMI, Reg, RegAGPRClass, TII, SRI))
return false;
SmallVector<MachineOperand *, 8> DefReachingUses;
findReachingUses(&DefMI, DAG.LIS, DefReachingUses);
@@ -2377,7 +2366,7 @@ bool RewriteMFMAFormStage::isRecolorSafe(
RewriteCandsSet)) {
continue;
}
- if (!useAcceptsAGPR(UseMO, TII, SRI))
+ if (!useAcceptsAGPR(UseMO, RegAGPRClass, TII, SRI))
return false;
}
}
@@ -2815,7 +2804,11 @@ bool RewriteMFMAFormStage::rewrite(
if (TII->isMAI(*UserMI))
CanReadAGPR = RewriteCandsSet.contains(UserMI);
else
- CanReadAGPR = DstRecolorSafe && useAcceptsAGPR(RUOp, TII, SRI);
+ CanReadAGPR =
+ DstRecolorSafe &&
+ useAcceptsAGPR(
+ RUOp, SRI->getEquivalentAGPRClass(DAG.MRI.getRegClass(DstReg)),
+ TII, SRI);
if (!CanReadAGPR &&
find(DstReachingUseCopies, RUOp) == DstReachingUseCopies.end())
DstReachingUseCopies.push_back(RUOp);
diff --git a/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir b/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir
index de9c11753db37..4e0206e628509 100644
--- a/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir
+++ b/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir
@@ -19,8 +19,7 @@ body: |
; CHECK-NEXT: successors: %bb.1(0x80000000)
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: %p0:vreg_1024 = IMPLICIT_DEF
- ; CHECK-NEXT: %m3:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = COPY %m3
+ ; CHECK-NEXT: %m3:areg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: %p1:vreg_1024 = IMPLICIT_DEF
; CHECK-NEXT: %p2:vreg_1024 = IMPLICIT_DEF
; CHECK-NEXT: %p3:vreg_1024 = IMPLICIT_DEF
@@ -33,7 +32,7 @@ body: |
; CHECK-NEXT: bb.1:
; CHECK-NEXT: successors: %bb.1(0x7c000000), %bb.2(0x04000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: %m0:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, [[COPY]], 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %m0:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m3, 0, 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: %m1:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m0, 0, 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: %m2:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m1, 0, 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: S_CBRANCH_SCC1 %bb.2, implicit undef $scc
@@ -44,9 +43,7 @@ body: |
; CHECK-NEXT: %m4:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m2, 0, 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: %out1:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m4, 0, 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: %out2:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s1:vreg_128_align2, undef %s0:vreg_128_align2, %m4, 0, 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY %out1
- ; CHECK-NEXT: [[COPY2:%[0-9]+]]:vreg_128_align2 = COPY %out2
- ; CHECK-NEXT: S_ENDPGM 0, implicit [[COPY1]], implicit [[COPY2]]
+ ; CHECK-NEXT: S_ENDPGM 0, implicit %out1, implicit %out2
bb.0:
successors: %bb.1(0x80000000)
%p0:vreg_1024 = IMPLICIT_DEF
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-recolor-no-use.mir b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-recolor-no-use.mir
index 83cc85686fb15..3ed11ebb1fb5e 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-recolor-no-use.mir
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-unsafe-recolor-no-use.mir
@@ -43,8 +43,7 @@ body: |
; CHECK-NEXT: [[DEF10:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF12:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = COPY [[DEF12]]
+ ; CHECK-NEXT: [[DEF12:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.1:
; CHECK-NEXT: successors: %bb.2(0x80000000)
@@ -53,17 +52,17 @@ body: |
; CHECK-NEXT: [[DEF14:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF16:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF13]], [[DEF14]], [[COPY]], 4, 4, [[DEF15]].sub0, [[DEF16]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF13]], [[DEF14]], [[DEF12]], 4, 4, [[DEF15]].sub0, [[DEF16]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF13]], [[DEF14]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]], 4, 4, [[DEF15]].sub0, [[DEF16]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF13]], [[DEF14]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]], 4, 4, [[DEF15]].sub0, [[DEF16]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF13]], [[DEF14]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]], 4, 4, [[DEF15]].sub0, [[DEF16]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]]
- ; CHECK-NEXT: [[COPY1:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[DEF11]].sub0, [[DEF16]], implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]]
+ ; CHECK-NEXT: [[COPY:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[DEF11]].sub0, [[DEF16]], implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF17:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
- ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF11]], [[DEF12]], [[COPY1]]
+ ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF11]], [[DEF12]], [[COPY]]
; CHECK-NEXT: S_NOP 0, implicit %12, implicit %13
; CHECK-NEXT: S_ENDPGM 0
bb.0:
diff --git a/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_copies.mir b/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_copies.mir
index 845ab48af9ab2..f7126dc1ab974 100644
--- a/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_copies.mir
+++ b/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_copies.mir
@@ -1791,10 +1791,13 @@ body: |
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: undef [[V_ADD_U32_e32_:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[DEF13]].sub1, [[DEF12]], implicit $exec
+ ; CHECK-NEXT: undef [[V_ADD_U32_e32_:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[DEF13]].sub0, [[DEF12]], implicit $exec
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = COPY [[V_ADD_U32_e32_]]
+ ; CHECK-NEXT: undef [[V_ADD_U32_e32_1:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[DEF13]].sub1, [[DEF12]], implicit $exec
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:areg_128_align2 = COPY [[V_ADD_U32_e32_1]]
; CHECK-NEXT: $scc = IMPLICIT_DEF
; CHECK-NEXT: S_CBRANCH_SCC1 %bb.4, implicit killed $scc
; CHECK-NEXT: {{ $}}
@@ -1804,28 +1807,27 @@ body: |
; CHECK-NEXT: [[DEF17:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF18:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF19:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: undef [[V_ADD_U32_e32_1:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[DEF13]].sub0, [[DEF12]], implicit $exec
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[V_ADD_U32_e32_1]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_ADD_U32_e32_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_1:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_2:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_4:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_5:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_6:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_7:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_8:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_9:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_10:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_11:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[COPY]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_4:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_5:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_6:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_7:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_8:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_9:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_10:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_11:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: KILL [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_1]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_2]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_4]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_5]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_6]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_7]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_8]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_9]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_10]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_11]]
+ ; CHECK-NEXT: KILL [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_4]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_5]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_6]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_7]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_8]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_9]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_10]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_11]]
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.3:
; CHECK-NEXT: successors: %bb.5(0x40000000), %bb.4(0x40000000)
@@ -1836,17 +1838,20 @@ body: |
; CHECK-NEXT: bb.4:
; CHECK-NEXT: successors: %bb.5(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: dead undef [[V_ADD_U32_e32_2:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[V_ADD_U32_e32_]].sub1, [[DEF12]], implicit $exec
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:vreg_128_align2 = COPY [[COPY1]]
+ ; CHECK-NEXT: dead undef [[V_ADD_U32_e32_2:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[COPY2]].sub1, [[DEF12]], implicit $exec
; CHECK-NEXT: S_BRANCH %bb.5
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.5:
; CHECK-NEXT: successors: %bb.6(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: undef [[V_ADD_U32_e32_3:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[V_ADD_U32_e32_]].sub0, [[DEF12]], implicit $exec
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[COPY1]]
+ ; CHECK-NEXT: undef [[V_ADD_U32_e32_3:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[COPY3]].sub0, [[DEF12]], implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.6:
+ ; CHECK-NEXT: [[COPY4:%[0-9]+]]:vreg_128_align2 = COPY [[COPY1]]
; CHECK-NEXT: [[DEF20:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
- ; CHECK-NEXT: KILL [[DEF20]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF11]], [[DEF13]], [[DEF14]], [[DEF15]], [[DEF16]], [[V_ADD_U32_e32_]], [[V_ADD_U32_e32_3]]
+ ; CHECK-NEXT: KILL [[DEF20]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF11]], [[DEF13]], [[DEF14]], [[DEF15]], [[DEF16]], [[COPY4]], [[V_ADD_U32_e32_3]]
; CHECK-NEXT: S_NOP 0, implicit %12, implicit %13
; CHECK-NEXT: S_ENDPGM 0
bb.0:
@@ -2469,8 +2474,7 @@ body: |
; CHECK-NEXT: [[DEF10:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = COPY [[DS_READ_B128_gfx9_]]
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: $scc = IMPLICIT_DEF
@@ -2482,34 +2486,27 @@ body: |
; CHECK-NEXT: [[DEF14:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF15:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_4:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_5:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[COPY]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_4:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_5:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF16]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_5]]
- ; CHECK-NEXT: [[COPY2:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]]
- ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]]
- ; CHECK-NEXT: [[COPY4:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_4]]
- ; CHECK-NEXT: [[COPY5:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]]
- ; CHECK-NEXT: [[COPY6:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]]
- ; CHECK-NEXT: KILL [[COPY2]], [[COPY5]], [[COPY3]], [[COPY6]], [[COPY4]], [[COPY1]]
+ ; CHECK-NEXT: KILL [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_4]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_5]]
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.3:
- ; CHECK-NEXT: [[COPY7:%[0-9]+]]:vreg_128_align2 = COPY [[COPY]]
- ; CHECK-NEXT: DS_WRITE_B128_gfx9 [[DEF11]], [[COPY7]], 0, 0, implicit $exec
+ ; CHECK-NEXT: DS_WRITE_B128_gfx9 [[DEF11]], [[DS_READ_B128_gfx9_]], 0, 0, implicit $exec
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF17:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
- ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF12]], [[DEF13]], [[COPY7]]
+ ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF12]], [[DEF13]], [[DS_READ_B128_gfx9_]]
; CHECK-NEXT: S_NOP 0, implicit %12, implicit %13
; CHECK-NEXT: S_ENDPGM 0
bb.0:
@@ -2626,20 +2623,14 @@ body: |
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_5]]
- ; CHECK-NEXT: [[COPY2:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]]
- ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]]
- ; CHECK-NEXT: [[COPY4:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_4]]
- ; CHECK-NEXT: [[COPY5:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]]
- ; CHECK-NEXT: [[COPY6:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]]
- ; CHECK-NEXT: KILL [[COPY2]], [[COPY5]], [[COPY3]], [[COPY6]], [[COPY4]], [[COPY1]]
+ ; CHECK-NEXT: KILL [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_4]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_5]]
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.3:
- ; CHECK-NEXT: [[COPY7:%[0-9]+]]:vreg_128_align2 = COPY [[COPY]]
- ; CHECK-NEXT: [[V_ADD_U32_e32_:%[0-9]+]]:vgpr_32 = V_ADD_U32_e32 [[DEF13]].sub1, [[COPY7]].sub0, implicit $exec
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY [[COPY]]
+ ; CHECK-NEXT: [[V_ADD_U32_e32_:%[0-9]+]]:vgpr_32 = V_ADD_U32_e32 [[DEF13]].sub1, [[COPY1]].sub0, implicit $exec
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF17:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
- ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF12]], [[DEF13]], [[COPY7]], [[V_ADD_U32_e32_]]
+ ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF12]], [[DEF13]], [[COPY1]], [[V_ADD_U32_e32_]]
; CHECK-NEXT: S_NOP 0, implicit %12, implicit %13
; CHECK-NEXT: S_ENDPGM 0
bb.0:
@@ -2960,13 +2951,13 @@ body: |
; CHECK-NEXT: [[DEF10:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 128, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 128, 0, implicit $exec
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: $scc = IMPLICIT_DEF
; CHECK-NEXT: S_CBRANCH_SCC1 %bb.2, implicit killed $scc
; CHECK-NEXT: {{ $}}
@@ -2976,10 +2967,10 @@ body: |
; CHECK-NEXT: [[DEF17:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF18:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF19:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: DS_WRITE_B128_gfx9 [[DEF11]], [[DS_READ_B128_gfx9_1]], 0, 0, implicit $exec
@@ -3069,13 +3060,13 @@ body: |
; CHECK-NEXT: [[DEF10:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 128, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 128, 0, implicit $exec
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: $scc = IMPLICIT_DEF
; CHECK-NEXT: S_CBRANCH_SCC1 %bb.2, implicit killed $scc
; CHECK-NEXT: {{ $}}
@@ -3085,10 +3076,10 @@ body: |
; CHECK-NEXT: [[DEF17:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF18:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF19:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: DS_WRITE_B128_gfx9 [[DEF11]], [[DS_READ_B128_gfx9_]], 0, 0, implicit $exec
@@ -3180,25 +3171,25 @@ body: |
; CHECK-NEXT: [[DEF10:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: $scc = IMPLICIT_DEF
; CHECK-NEXT: S_CBRANCH_SCC1 %bb.2, implicit killed $scc
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.1:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 256, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 256, 0, implicit $exec
; CHECK-NEXT: S_BRANCH %bb.3
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 512, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 512, 0, implicit $exec
; CHECK-NEXT: S_BRANCH %bb.3
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.3:
@@ -3213,10 +3204,10 @@ body: |
; CHECK-NEXT: [[DEF17:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF18:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF19:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_1]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_1]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.5:
; CHECK-NEXT: DS_WRITE_B128_gfx9 [[DEF11]], [[DS_READ_B128_gfx9_]], 128, 0, implicit $exec
@@ -3316,25 +3307,25 @@ body: |
; CHECK-NEXT: [[DEF10:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: $scc = IMPLICIT_DEF
; CHECK-NEXT: S_CBRANCH_SCC1 %bb.2, implicit killed $scc
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.1:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 256, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 256, 0, implicit $exec
; CHECK-NEXT: S_BRANCH %bb.3
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 512, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 512, 0, implicit $exec
; CHECK-NEXT: S_BRANCH %bb.3
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.3:
@@ -3349,10 +3340,10 @@ body: |
; CHECK-NEXT: [[DEF17:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF18:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF19:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_1]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_1]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.5:
; CHECK-NEXT: DS_WRITE_B128_gfx9 [[DEF11]], [[DS_READ_B128_gfx9_1]], 128, 0, implicit $exec
@@ -3613,12 +3604,8 @@ body: |
; CHECK-NEXT: bb.5:
; CHECK-NEXT: DS_WRITE_B128_gfx9 [[DEF11]], [[DS_READ_B128_gfx9_]], 0, 0, implicit $exec
; CHECK-NEXT: SCHED_BARRIER 0
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]]
- ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]]
- ; CHECK-NEXT: [[COPY2:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]]
- ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]]
; CHECK-NEXT: [[DEF17:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
- ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF12]], [[DEF13]], [[COPY]], [[COPY2]], [[COPY1]], [[COPY3]]
+ ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF12]], [[DEF13]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]]
; CHECK-NEXT: S_NOP 0, implicit %12, implicit %13
; CHECK-NEXT: S_ENDPGM 0
bb.0:
@@ -4004,13 +3991,13 @@ body: |
; CHECK-NEXT: [[DEF10:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 256, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 256, 0, implicit $exec
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: $scc = IMPLICIT_DEF
; CHECK-NEXT: S_CBRANCH_SCC1 %bb.3, implicit killed $scc
; CHECK-NEXT: {{ $}}
@@ -4020,10 +4007,10 @@ body: |
; CHECK-NEXT: [[DEF17:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF18:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF19:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.4(0x40000000), %bb.3(0x40000000)
@@ -4138,13 +4125,13 @@ body: |
; CHECK-NEXT: [[DEF10:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 256, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 256, 0, implicit $exec
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: $scc = IMPLICIT_DEF
; CHECK-NEXT: S_CBRANCH_SCC1 %bb.3, implicit killed $scc
; CHECK-NEXT: {{ $}}
@@ -4154,10 +4141,10 @@ body: |
; CHECK-NEXT: [[DEF17:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF18:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF19:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.4(0x40000000), %bb.3(0x40000000)
@@ -4280,25 +4267,25 @@ body: |
; CHECK-NEXT: [[DEF10:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: $scc = IMPLICIT_DEF
; CHECK-NEXT: S_CBRANCH_SCC1 %bb.2, implicit killed $scc
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.1:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 256, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 256, 0, implicit $exec
; CHECK-NEXT: S_BRANCH %bb.3
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 512, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 512, 0, implicit $exec
; CHECK-NEXT: S_BRANCH %bb.3
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.3:
@@ -4313,10 +4300,10 @@ body: |
; CHECK-NEXT: [[DEF17:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF18:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF19:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_1]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_1]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.5:
; CHECK-NEXT: successors: %bb.7(0x40000000), %bb.6(0x40000000)
@@ -4445,25 +4432,25 @@ body: |
; CHECK-NEXT: [[DEF10:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 0, 0, implicit $exec
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: $scc = IMPLICIT_DEF
; CHECK-NEXT: S_CBRANCH_SCC1 %bb.2, implicit killed $scc
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.1:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 256, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 256, 0, implicit $exec
; CHECK-NEXT: S_BRANCH %bb.3
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:vreg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 512, 0, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_1:%[0-9]+]]:areg_128_align2 = DS_READ_B128_gfx9 [[DEF11]], 512, 0, implicit $exec
; CHECK-NEXT: S_BRANCH %bb.3
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.3:
@@ -4478,10 +4465,10 @@ body: |
; CHECK-NEXT: [[DEF17:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF18:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF19:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_1]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF15:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DS_READ_B128_gfx9_1]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF15:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF14]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF16:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF15]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DS_READ_B128_gfx9_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[DEF16]], 4, 4, [[DEF19]].sub0, [[DEF11]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.5:
; CHECK-NEXT: successors: %bb.7(0x40000000), %bb.6(0x40000000)
@@ -5490,7 +5477,7 @@ body: |
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF13:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.1:
; CHECK-NEXT: successors: %bb.2(0x80000000)
@@ -5498,17 +5485,18 @@ body: |
; CHECK-NEXT: [[DEF14:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF15:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF16:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[DEF13]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_1:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_2:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_1]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_2]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[DEF13]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_1]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF14]], [[DEF15]], [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_2]], 4, 4, [[DEF16]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
- ; CHECK-NEXT: undef [[V_ADD_U32_e32_:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_]].sub0, [[DEF12]], implicit $exec
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64_3]]
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_]]
+ ; CHECK-NEXT: undef [[V_ADD_U32_e32_:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[COPY]].sub0, [[DEF12]], implicit $exec
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY [[V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64_3]]
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF17:%[0-9]+]]:vreg_1024 = IMPLICIT_DEF
- ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF11]], [[DEF13]], [[V_ADD_U32_e32_]], [[COPY]]
+ ; CHECK-NEXT: KILL [[DEF17]], [[DEF]], [[DEF1]], [[DEF2]], [[DEF3]], [[DEF4]], [[DEF5]], [[DEF6]], [[DEF7]], [[DEF8]], [[DEF9]], [[DEF10]], [[DEF11]], [[DEF13]], [[V_ADD_U32_e32_]], [[COPY1]]
; CHECK-NEXT: S_NOP 0, implicit %12, implicit %13
; CHECK-NEXT: S_ENDPGM 0
bb.0:
diff --git a/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_cost.mir b/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_cost.mir
index 32375684cbf99..bbd736754456d 100644
--- a/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_cost.mir
+++ b/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_cost.mir
@@ -242,7 +242,7 @@ body: |
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF13:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: dead [[DEF15:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: dead [[DEF16:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: dead undef [[V_ADD_U32_e32_:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[DEF13]].sub1, [[DEF12]], implicit $exec
@@ -253,12 +253,14 @@ body: |
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: undef [[V_ADD_U32_e32_1:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[DEF13]].sub0, [[DEF12]], implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = COPY [[V_ADD_U32_e32_1]]
; CHECK-NEXT: S_BRANCH %bb.3
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: undef [[V_ADD_U32_e32_1:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[DEF13]].sub1, [[DEF12]], implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = COPY [[V_ADD_U32_e32_1]]
; CHECK-NEXT: S_BRANCH %bb.3
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.3:
@@ -273,7 +275,7 @@ body: |
; CHECK-NEXT: [[DEF17:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF18:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF19:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF14:%[0-9]+]]:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[V_ADD_U32_e32_1]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF14:%[0-9]+]]:areg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[COPY]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.5:
; CHECK-NEXT: successors: %bb.7(0x40000000), %bb.6(0x40000000)
@@ -284,15 +286,17 @@ body: |
; CHECK-NEXT: bb.6:
; CHECK-NEXT: successors: %bb.8(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: undef [[V_ADD_U32_e32_2:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[DEF14]].sub1, [[DEF12]], implicit $exec
- ; CHECK-NEXT: [[V_ADD_U32_e32_2:%[0-9]+]].sub1:vreg_128_align2 = V_ADD_U32_e32 [[DEF14]].sub0, [[DEF12]], implicit $exec
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY [[DEF14]]
+ ; CHECK-NEXT: undef [[V_ADD_U32_e32_2:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[COPY1]].sub1, [[DEF12]], implicit $exec
+ ; CHECK-NEXT: [[V_ADD_U32_e32_2:%[0-9]+]].sub1:vreg_128_align2 = V_ADD_U32_e32 [[COPY1]].sub0, [[DEF12]], implicit $exec
; CHECK-NEXT: S_BRANCH %bb.8
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.7:
; CHECK-NEXT: successors: %bb.8(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: undef [[V_ADD_U32_e32_2:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[DEF14]].sub0, [[DEF12]], implicit $exec
- ; CHECK-NEXT: [[V_ADD_U32_e32_2:%[0-9]+]].sub1:vreg_128_align2 = V_ADD_U32_e32 [[DEF14]].sub1, [[DEF12]], implicit $exec
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:vreg_128_align2 = COPY [[DEF14]]
+ ; CHECK-NEXT: undef [[V_ADD_U32_e32_2:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[COPY2]].sub0, [[DEF12]], implicit $exec
+ ; CHECK-NEXT: [[V_ADD_U32_e32_2:%[0-9]+]].sub1:vreg_128_align2 = V_ADD_U32_e32 [[COPY2]].sub1, [[DEF12]], implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.8:
; CHECK-NEXT: [[V_ADD_U32_e32_3:%[0-9]+]]:vgpr_32 = V_ADD_U32_e32 [[DEF13]].sub1, [[V_ADD_U32_e32_1]].sub0, implicit $exec
diff --git a/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_diff_types.mir b/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_diff_types.mir
index 7528b855e23e3..e4a7e1d3e2e24 100644
--- a/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_diff_types.mir
+++ b/llvm/test/CodeGen/AMDGPU/sched_mfma_rewrite_diff_types.mir
@@ -42,7 +42,7 @@ body: |
; CHECK-NEXT: [[DEF9:%[0-9]+]]:vreg_128 = IMPLICIT_DEF
; CHECK-NEXT: S_NOP 0, implicit-def %12
; CHECK-NEXT: S_NOP 0, implicit-def %13
- ; CHECK-NEXT: [[DEF10:%[0-9]+]]:vreg_1024_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF10:%[0-9]+]]:areg_1024_align2 = IMPLICIT_DEF
; CHECK-NEXT: SCHED_BARRIER 0
; CHECK-NEXT: [[DEF11:%[0-9]+]]:vreg_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF12:%[0-9]+]]:vgpr_32 = IMPLICIT_DEF
@@ -58,12 +58,14 @@ body: |
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: undef [[V_ADD_U32_e32_1:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[DEF13]].sub0, [[DEF12]], implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = COPY [[V_ADD_U32_e32_1]]
; CHECK-NEXT: S_BRANCH %bb.3
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: undef [[V_ADD_U32_e32_1:%[0-9]+]].sub0:vreg_128_align2 = V_ADD_U32_e32 [[DEF13]].sub1, [[DEF12]], implicit $exec
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = COPY [[V_ADD_U32_e32_1]]
; CHECK-NEXT: S_BRANCH %bb.3
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.3:
@@ -78,7 +80,7 @@ body: |
; CHECK-NEXT: [[DEF17:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF18:%[0-9]+]]:av_128_align2 = IMPLICIT_DEF
; CHECK-NEXT: [[DEF19:%[0-9]+]]:vreg_64_align2 = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF10:%[0-9]+]].sub0_sub1_sub2_sub3:vreg_1024_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 [[DEF17]], [[DEF18]], [[V_ADD_U32_e32_1]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[DEF10:%[0-9]+]].sub0_sub1_sub2_sub3:areg_1024_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_e64 [[DEF17]], [[DEF18]], [[COPY]], 4, 4, [[DEF19]].sub0, [[DEF12]], 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.5:
; CHECK-NEXT: successors: %bb.7(0x40000000), %bb.6(0x40000000)
>From 8976dc1968faca2272d8b8575d35d627a585d8ad Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Tue, 8 Sep 2026 11:48:28 +0000
Subject: [PATCH 7/9] [AMDGPU] Simplify MFMA recolor AGPR-constraint checks
Use getRegClassConstraintEffectForVReg so a tied def/use is constrained
on both sides, drop the dead RegAGPRClass null checks (the class always
comes from getEquivalentAGPRClass, which never returns null), and fold
the CanReadAGPR def/use branch into a single expression.
Co-Authored-By: Claude Opus 4.7 <noreply at anthropic.com>
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 25 ++++++++-------------
1 file changed, 9 insertions(+), 16 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index 166500c02630c..11627efe31c36 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -2323,20 +2323,16 @@ isRewriteCandidateMAI(const MachineInstr *MI, const SIInstrInfo *TII,
static bool canWriteAGPR(const MachineInstr *MI, Register Reg,
const TargetRegisterClass *RegAGPRClass,
const SIInstrInfo *TII, const SIRegisterInfo *SRI) {
- if (MI->getDesc().getNumDefs() == 0 || !RegAGPRClass)
+ if (MI->getDesc().getNumDefs() == 0)
return false;
- int DefOpIdx =
- MI->findRegisterDefOperandIdx(Reg, /*TRI=*/nullptr, false, false);
- return DefOpIdx >= 0 && MI->getRegClassConstraintEffect(
- DefOpIdx, RegAGPRClass, TII, SRI) != nullptr;
+ return MI->getRegClassConstraintEffectForVReg(Reg, RegAGPRClass, TII, SRI) !=
+ nullptr;
}
static bool useAcceptsAGPR(const MachineOperand *Use,
const TargetRegisterClass *RegAGPRClass,
const SIInstrInfo *TII, const SIRegisterInfo *SRI) {
const MachineInstr *UseMI = Use->getParent();
- if (!RegAGPRClass)
- return false;
return UseMI->getRegClassConstraintEffect(Use->getOperandNo(), RegAGPRClass,
TII, SRI) != nullptr;
}
@@ -2791,6 +2787,8 @@ bool RewriteMFMAFormStage::rewrite(
bool DstRecolorSafe =
!DstAlreadyRedef &&
isRecolorSafe(DstReg, DstReachingUses, RewriteCandsSet, /*IsDst=*/true);
+ const TargetRegisterClass *DstAGPRClass =
+ SRI->getEquivalentAGPRClass(DAG.MRI.getRegClass(DstReg));
for (MachineOperand *RUOp : DstReachingUses) {
MachineInstr *UserMI = RUOp->getParent();
// Decide whether this reaching use can read the dst's AGPR form directly
@@ -2800,15 +2798,10 @@ bool RewriteMFMAFormStage::rewrite(
// to AGPR (DstRecolorSafe) and its operand accepts an AGPR. When the
// dst is unsafe, its original reg stays VGPR, so every non-MFMA user
// must go through a bridge copy.
- bool CanReadAGPR;
- if (TII->isMAI(*UserMI))
- CanReadAGPR = RewriteCandsSet.contains(UserMI);
- else
- CanReadAGPR =
- DstRecolorSafe &&
- useAcceptsAGPR(
- RUOp, SRI->getEquivalentAGPRClass(DAG.MRI.getRegClass(DstReg)),
- TII, SRI);
+ bool CanReadAGPR =
+ TII->isMAI(*UserMI)
+ ? RewriteCandsSet.contains(UserMI)
+ : DstRecolorSafe && useAcceptsAGPR(RUOp, DstAGPRClass, TII, SRI);
if (!CanReadAGPR &&
find(DstReachingUseCopies, RUOp) == DstReachingUseCopies.end())
DstReachingUseCopies.push_back(RUOp);
>From c0dea558be3635071b3699f64dad2d4d1af54d21 Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Thu, 10 Sep 2026 01:40:40 +0000
Subject: [PATCH 8/9] [AMDGPU] Rename useAcceptsAGPR to userAcceptsAGPR and key
on user instruction
Pass the user MachineInstr and register instead of an operand, using
getRegClassConstraintEffectForVReg so the constraint check iterates the
instruction's operands. Addresses review feedback.
Co-Authored-By: Claude Opus 4.7 <noreply at anthropic.com>
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 16 ++++++++--------
1 file changed, 8 insertions(+), 8 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index 69e80ab9d734d..05067baf6edf4 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -2366,12 +2366,11 @@ static bool canWriteAGPR(const MachineInstr *MI, Register Reg,
nullptr;
}
-static bool useAcceptsAGPR(const MachineOperand *Use,
- const TargetRegisterClass *RegAGPRClass,
- const SIInstrInfo *TII, const SIRegisterInfo *SRI) {
- const MachineInstr *UseMI = Use->getParent();
- return UseMI->getRegClassConstraintEffect(Use->getOperandNo(), RegAGPRClass,
- TII, SRI) != nullptr;
+static bool userAcceptsAGPR(const MachineInstr *UserMI, Register Reg,
+ const TargetRegisterClass *RegAGPRClass,
+ const SIInstrInfo *TII, const SIRegisterInfo *SRI) {
+ return UserMI->getRegClassConstraintEffectForVReg(Reg, RegAGPRClass, TII,
+ SRI) != nullptr;
}
bool RewriteMFMAFormStage::isRecolorSafe(
@@ -2399,7 +2398,7 @@ bool RewriteMFMAFormStage::isRecolorSafe(
RewriteCandsSet)) {
continue;
}
- if (!useAcceptsAGPR(UseMO, RegAGPRClass, TII, SRI))
+ if (!userAcceptsAGPR(UseMO->getParent(), Reg, RegAGPRClass, TII, SRI))
return false;
}
}
@@ -2838,7 +2837,8 @@ bool RewriteMFMAFormStage::rewrite(
bool CanReadAGPR =
TII->isMAI(*UserMI)
? RewriteCandsSet.contains(UserMI)
- : DstRecolorSafe && useAcceptsAGPR(RUOp, DstAGPRClass, TII, SRI);
+ : DstRecolorSafe &&
+ userAcceptsAGPR(UserMI, DstReg, DstAGPRClass, TII, SRI);
if (!CanReadAGPR &&
find(DstReachingUseCopies, RUOp) == DstReachingUseCopies.end())
DstReachingUseCopies.push_back(RUOp);
>From 356beabcd33540bb1bef4a95ab28577bb8990623 Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Mon, 14 Sep 2026 09:38:57 +0000
Subject: [PATCH 9/9] [AMDGPU][NFC] Simplify RewriteSet population in
initHeuristics
Fold the candidate check into the insert instead of using a continue,
making the loop body a single conditional insert.
Co-Authored-By: Claude Opus 4.7 <noreply at anthropic.com>
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 5 ++---
1 file changed, 2 insertions(+), 3 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index 05067baf6edf4..6d5c9f8c3990d 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -2447,9 +2447,8 @@ bool RewriteMFMAFormStage::initHeuristics(
SmallPtrSet<MachineInstr *, 16> RewriteSet;
for (MachineBasicBlock &MBB : MF) {
for (MachineInstr &MI : MBB) {
- if (!isRewriteCandidate(&MI))
- continue;
- RewriteSet.insert(&MI);
+ if (isRewriteCandidate(&MI))
+ RewriteSet.insert(&MI);
}
}
More information about the llvm-commits
mailing list