[llvm] [AMDGPU] Add block carried latency to CoExecSched (PR #192324)
Austin Kerbow via llvm-commits
llvm-commits at lists.llvm.org
Fri Sep 11 11:09:56 PDT 2026
https://github.com/kerbowa updated https://github.com/llvm/llvm-project/pull/192324
>From a02640eb59d0b48355b192420f1522edb43133c9 Mon Sep 17 00:00:00 2001
From: Jeffrey Byrnes <Jeffrey.Byrnes at amd.com>
Date: Wed, 18 Mar 2026 12:16:52 -0700
Subject: [PATCH 01/10] [AMDGPU] Add block carried latency to CoExecSched
Change-Id: Ib04e40e57d38e127d6c5452d1719e32dacef2ade
---
.../AMDGPU/AMDGPUCoExecSchedStrategy.cpp | 260 +++--
.../Target/AMDGPU/AMDGPUCoExecSchedStrategy.h | 28 +-
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 36 -
llvm/lib/Target/AMDGPU/GCNSchedStrategy.h | 4 -
llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll | 904 +++++++++++++++++-
5 files changed, 1112 insertions(+), 120 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
index 1f3905a0465e4..9f23ddb7cf6bb 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
@@ -544,12 +544,15 @@ CandidateHeuristics::getHWUIFromFlavor(InstructionFlavor Flavor) {
return nullptr;
}
-unsigned CandidateHeuristics::getHWUICyclesForInst(SUnit *SU) {
+unsigned CandidateHeuristics::getHWUICyclesForSU(SUnit *SU) {
assert(SchedModel && SchedModel->hasInstrSchedModel());
MachineInstr *MI = SU->getInstr();
if (SII->isDS(*MI))
return SchedModel->computeInstrLatency(MI);
+ // Loads and stores are not pipelined
+ if (MI->mayLoadOrStore())
+ return SchedModel->computeInstrLatency(MI, false);
unsigned ReleaseAtCycle = 0;
const MCSchedClassDesc *SC = DAG->getSchedClass(SU);
for (TargetSchedModel::ProcResIter PI = SchedModel->getWriteProcResBegin(SC),
@@ -564,7 +567,23 @@ void CandidateHeuristics::updateForScheduling(SUnit *SU) {
HardwareUnitInfo *HWUI =
getHWUIFromFlavor(classifyFlavor(*SU->getInstr(), *SII));
assert(HWUI);
- HWUI->markScheduled(SU, getHWUICyclesForInst(SU));
+ HWUI->markScheduled(SU, getHWUICyclesForSU(SU));
+}
+
+unsigned CandidateHeuristics::getHWUICyclesForMI(MachineInstr *MI) {
+ assert(SchedModel && SchedModel->hasInstrSchedModel());
+ // Loads and stores are not pipelined
+ if (MI->mayLoadOrStore())
+ return SchedModel->computeInstrLatency(MI, false);
+
+ unsigned ReleaseAtCycle = 0;
+ const MCSchedClassDesc *SC = SchedModel->resolveSchedClass(MI);
+ for (TargetSchedModel::ProcResIter PI = SchedModel->getWriteProcResBegin(SC),
+ PE = SchedModel->getWriteProcResEnd(SC);
+ PI != PE; ++PI) {
+ ReleaseAtCycle = std::max(ReleaseAtCycle, (unsigned)PI->ReleaseAtCycle);
+ }
+ return ReleaseAtCycle;
}
void CandidateHeuristics::initialize(ScheduleDAGMI *SchedDAG,
@@ -589,16 +608,77 @@ void CandidateHeuristics::initialize(ScheduleDAGMI *SchedDAG,
HWUInfo[(int)InstructionFlavor::TRANS].setProducesCoexecWindow(true);
HWUInfo[(int)InstructionFlavor::DS].setBufferSize(DefaultBufferSizes::DS);
- collectHWUIPressure();
+ collectRegionSummary();
+}
+
+unsigned CandidateHeuristics::getCarriedLatency(SUnit *SU) {
+ MachineInstr *MI = SU->getInstr();
+ unsigned CarriedLatency = 0;
+ for (auto &Op : MI->operands()) {
+ if (!Op.isReg())
+ continue;
+ if (!Op.isUse())
+ continue;
+ auto Reg = Op.getReg();
+ if (!Reg.isVirtual())
+ continue;
+
+ for (auto &Def : DAG->MRI.def_instructions(Reg)) {
+ // We don't have the proper modelling to accurately measure all carried
+ // latency. Just try to measure carried latency for long latency loads to
+ // avoid long stalls.
+ if (!Def.mayLoad())
+ continue;
+
+ unsigned Latency = getHWUICyclesForMI(&Def);
+
+ // Load is carried across block
+ if (Def.getParent() != MI->getParent()) {
+ bool FoundUseInDefBlock = false;
+ for (auto &Use : DAG->MRI.use_nodbg_instructions(Reg)) {
+ if (Use.getParent() != Def.getParent())
+ continue;
+
+ SlotIndex DefIdx = DAG->getLIS()->getInstructionIndex(Def);
+ SlotIndex UseIdx = DAG->getLIS()->getInstructionIndex(Use);
+ // We have a use of this load in the def block that occurs after the
+ // load. In this case we must wait for the load in the def block, and
+ // we do not have any carried latency from this load.
+ if (SlotIndex::isEarlierInstr(DefIdx, UseIdx)) {
+ FoundUseInDefBlock = true;
+ break;
+ }
+ }
+ if (!FoundUseInDefBlock)
+ CarriedLatency = std::max(Latency, CarriedLatency);
+
+ continue;
+ }
+
+ assert(Def.getParent() == MI->getParent());
+ // Load is in the same block
+ SlotIndex LoadIdx = DAG->getLIS()->getInstructionIndex(Def);
+ SlotIndex UseIdx = DAG->getLIS()->getInstructionIndex(*MI);
+ // The load occurs after this use -- the latency is carried across loop
+ // backedge.
+ if (SlotIndex::isEarlierInstr(UseIdx, LoadIdx))
+ CarriedLatency = std::max(Latency, CarriedLatency);
+ }
+ }
+ return CarriedLatency;
}
-void CandidateHeuristics::collectHWUIPressure() {
+void CandidateHeuristics::collectRegionSummary() {
if (!SchedModel || !SchedModel->hasInstrSchedModel())
return;
for (auto &SU : DAG->SUnits) {
- const InstructionFlavor Flavor = classifyFlavor(*SU.getInstr(), *SII);
- HWUInfo[(int)(Flavor)].insert(&SU, getHWUICyclesForInst(&SU));
+ MachineInstr *MI = SU.getInstr();
+ const InstructionFlavor Flavor = classifyFlavor(*MI, *SII);
+ HWUInfo[(int)(Flavor)].insert(&SU, getHWUICyclesForSU(&SU));
+ unsigned CarriedLatency = getCarriedLatency(&SU);
+ if (CarriedLatency)
+ CarriedLatencies[MI] = CarriedLatency;
}
for (auto &HWUI : HWUInfo)
@@ -645,6 +725,111 @@ void CandidateHeuristics::sortHWUIResources() {
});
}
+unsigned CandidateHeuristics::getStructuralStallCycles(SchedBoundary &Zone,
+ SUnit *SU) {
+ // Only implemented for top-down scheduling currently.
+ if (!Zone.isTop() || !SU)
+ return 0;
+
+ MachineInstr *MI = SU->getInstr();
+ unsigned CurrCycle = Zone.getCurrCycle();
+ unsigned Stall = 0;
+
+ // Query SchedModel for resource stalls (unbuffered resources).
+ if (SchedModel->hasInstrSchedModel() && SU->hasReservedResource) {
+ const MCSchedClassDesc *SC = DAG->getSchedClass(SU);
+ for (const MCWriteProcResEntry &PE :
+ make_range(SchedModel->getWriteProcResBegin(SC),
+ SchedModel->getWriteProcResEnd(SC))) {
+ unsigned NextAvail =
+ Zone.getNextResourceCycle(SC, PE.ProcResourceIdx, PE.ReleaseAtCycle,
+ PE.AcquireAtCycle)
+ .first;
+ if (NextAvail > CurrCycle)
+ Stall = std::max(Stall, NextAvail - CurrCycle);
+ }
+ }
+
+ // Query HazardRecognizer for sequence-dependent hazard penalties.
+ if (Zone.HazardRec && Zone.HazardRec->isEnabled()) {
+ auto *HR = static_cast<GCNHazardRecognizer *>(Zone.HazardRec.get());
+ Stall = std::max(Stall, HR->getHazardWaitStates(MI));
+ }
+
+ return Stall;
+}
+
+bool CandidateHeuristics::tryEffectiveStall(
+ GenericSchedulerBase::SchedCandidate &Cand,
+ GenericSchedulerBase::SchedCandidate &TryCand, SchedBoundary &Zone) {
+
+ auto getBufferFullStalls = [this, &Zone](SUnit *SU) -> unsigned {
+ InstructionFlavor Flavor = classifyFlavor(
+ *SU->getInstr(), *static_cast<const SIInstrInfo *>(DAG->TII));
+ HardwareUnitInfo *HWUI = getHWUIFromFlavor(Flavor);
+
+ // A BufferSize of 0 means "unlimited" buffer, thus we will never fill it.
+ if (HWUI->getBufferSize() == 0)
+ return 0;
+
+ // getBufferAvailableCycle assumes top-down scheduling.
+ assert(Zone.isTop());
+ unsigned CurrCycle = Zone.getCurrCycle();
+ unsigned BufferReadyCycle = HWUI->getBufferAvailableCycle(CurrCycle);
+ if (BufferReadyCycle <= CurrCycle)
+ return 0;
+
+ return BufferReadyCycle - CurrCycle;
+ };
+
+ // Treat structural and latency stalls as a single scheduling cost for the
+ // current cycle.
+ struct StallCosts {
+ unsigned Ready = 0;
+ unsigned Structural = 0;
+ unsigned Latency = 0;
+ unsigned Effective = 0;
+ unsigned Buffer = 0;
+ unsigned Carried = 0;
+ };
+
+ unsigned CurrCycle = Zone.getCurrCycle();
+ auto GetStallCosts = [&](SUnit *SU) {
+ unsigned ReadyCycle = Zone.isTop() ? SU->TopReadyCycle : SU->BotReadyCycle;
+ StallCosts Costs;
+ Costs.Ready = ReadyCycle > CurrCycle ? ReadyCycle - CurrCycle : 0;
+ Costs.Structural = getStructuralStallCycles(Zone, SU);
+ Costs.Latency = Zone.getLatencyStallCycles(SU);
+ Costs.Buffer = getBufferFullStalls(SU);
+ unsigned TryCarriedLatency =
+ CarriedLatencies.contains(TryCand.SU->getInstr())
+ ? CarriedLatencies[TryCand.SU->getInstr()]
+ : 0;
+ Costs.Carried =
+ TryCarriedLatency > CurrCycle ? TryCarriedLatency - CurrCycle : 0;
+ Costs.Effective = std::max({Costs.Ready, Costs.Structural, Costs.Latency,
+ Costs.Buffer, Costs.Carried});
+ return Costs;
+ };
+
+ StallCosts TryCosts = GetStallCosts(TryCand.SU);
+ StallCosts CandCosts = GetStallCosts(Cand.SU);
+
+ LLVM_DEBUG(if (TryCosts.Effective || CandCosts.Effective) {
+ dbgs() << "Effective stalls: try=" << TryCosts.Effective
+ << " (ready=" << TryCosts.Ready << ", struct=" << TryCosts.Structural
+ << ", lat=" << TryCosts.Latency << ", buffer=" << TryCosts.Buffer
+ << ", carried=" << TryCosts.Carried
+ << ") cand=" << CandCosts.Effective << " (ready=" << CandCosts.Ready
+ << ", struct=" << CandCosts.Structural
+ << ", lat=" << CandCosts.Latency << ", buffer=" << CandCosts.Buffer
+ << ", carried=" << CandCosts.Carried << ")\n";
+ });
+
+ return tryLess(TryCosts.Effective, CandCosts.Effective, TryCand, Cand,
+ AMDGPUCoExecSchedStrategy::Stall);
+}
+
bool CandidateHeuristics::tryCriticalResourceDependency(
GenericSchedulerBase::SchedCandidate &TryCand,
GenericSchedulerBase::SchedCandidate &Cand, SchedBoundary *Zone) const {
@@ -989,7 +1174,7 @@ bool AMDGPUCoExecSchedStrategy::tryCandidateCoexec(SchedCandidate &Cand,
if (SameBoundary) {
// Compare candidates by the stall they would introduce if
// scheduled in the current cycle.
- if (tryEffectiveStall(Cand, TryCand, *Zone))
+ if (Heurs.tryEffectiveStall(Cand, TryCand, *Zone))
return TryCand.Reason != NoCand;
Heurs.sortHWUIResources();
@@ -1052,67 +1237,6 @@ bool AMDGPUCoExecSchedStrategy::tryCandidateCoexec(SchedCandidate &Cand,
return false;
}
-bool AMDGPUCoExecSchedStrategy::tryEffectiveStall(SchedCandidate &Cand,
- SchedCandidate &TryCand,
- SchedBoundary &Zone) {
- auto getBufferFullStalls = [this, &Zone](SUnit *SU) -> unsigned {
- InstructionFlavor Flavor = classifyFlavor(
- *SU->getInstr(), *static_cast<const SIInstrInfo *>(DAG->TII));
- HardwareUnitInfo *HWUI = Heurs.getHWUIFromFlavor(Flavor);
-
- // A BufferSize of 0 is unlimited, so it has no FIFO scheduling cost.
- if (HWUI->getBufferSize() == 0)
- return 0;
-
- // getBufferAvailableCycle assumes top-down scheduling.
- assert(Zone.isTop());
- unsigned CurrCycle = Zone.getCurrCycle();
- unsigned BufferReadyCycle = HWUI->getBufferAvailableCycle(CurrCycle);
- if (BufferReadyCycle <= CurrCycle)
- return 0;
-
- return BufferReadyCycle - CurrCycle;
- };
-
- // Treat structural and latency stalls as a single scheduling cost for the
- // current cycle.
- struct StallCosts {
- unsigned Ready = 0;
- unsigned Structural = 0;
- unsigned Latency = 0;
- unsigned Effective = 0;
- unsigned Buffer = 0;
- };
-
- unsigned CurrCycle = Zone.getCurrCycle();
- auto GetStallCosts = [&](SUnit *SU) {
- unsigned ReadyCycle = Zone.isTop() ? SU->TopReadyCycle : SU->BotReadyCycle;
- StallCosts Costs;
- Costs.Ready = ReadyCycle > CurrCycle ? ReadyCycle - CurrCycle : 0;
- Costs.Structural = getStructuralStallCycles(Zone, SU);
- Costs.Latency = Zone.getLatencyStallCycles(SU);
- Costs.Buffer = getBufferFullStalls(SU);
- Costs.Effective =
- std::max({Costs.Ready, Costs.Structural, Costs.Latency, Costs.Buffer});
- return Costs;
- };
-
- StallCosts TryCosts = GetStallCosts(TryCand.SU);
- StallCosts CandCosts = GetStallCosts(Cand.SU);
-
- LLVM_DEBUG(if (TryCosts.Effective || CandCosts.Effective) {
- dbgs() << "Effective stalls: try=" << TryCosts.Effective
- << " (ready=" << TryCosts.Ready << ", struct=" << TryCosts.Structural
- << ", lat=" << TryCosts.Latency << ", buffer=" << TryCosts.Buffer
- << ") cand=" << CandCosts.Effective << " (ready=" << CandCosts.Ready
- << ", struct=" << CandCosts.Structural
- << ", lat=" << CandCosts.Latency << ", buffer=" << CandCosts.Buffer
- << ")\n";
- });
-
- return tryLess(TryCosts.Effective, CandCosts.Effective, TryCand, Cand, Stall);
-}
-
ScheduleDAGInstrs *
llvm::createGCNCoExecMachineScheduler(MachineSchedContext *C) {
LLVM_DEBUG(dbgs() << "AMDGPU coexec preRA scheduler selected for "
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.h b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.h
index 672c5e3c33b97..5685091fc09c9 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.h
@@ -203,13 +203,25 @@ class CandidateHeuristics {
const SIRegisterInfo *SRI;
const TargetSchedModel *SchedModel;
SmallVector<HardwareUnitInfo, 8> HWUInfo;
+ DenseMap<MachineInstr *, unsigned> CarriedLatencies;
- /// Walk over the region and collect total usage per HardwareUnit.
- void collectHWUIPressure();
+ /// Walk over the region and collect characteristics for the various
+ /// heuristics.
+ void collectRegionSummary();
/// Compute the blocking cycles for the appropriate HardwareUnit given an \p
- /// SU.
- unsigned getHWUICyclesForInst(SUnit *SU);
+ /// SU
+ unsigned getHWUICyclesForSU(SUnit *SU);
+ /// Compute the blocking cycles for the appropriate HardwareUnit given an \p
+ /// MI
+ unsigned getHWUICyclesForMI(MachineInstr *MI);
+
+ /// Estimate the block carried latency from loads for a given \p SU. This is
+ /// essentially global scheduling info that our local scheduling
+ /// infrastructure lacks the necessary infrastructure to accurately measure.
+ /// Thus, this method just attempts to find a reasonable upper bound for
+ /// carried load latency to avoid long stalls.
+ unsigned getCarriedLatency(SUnit *SU);
public:
CandidateHeuristics() = default;
@@ -229,6 +241,12 @@ class CandidateHeuristics {
/// keeping the critical HardwareUnit busy.
void sortHWUIResources();
+ unsigned getStructuralStallCycles(SchedBoundary &Zone, SUnit *SU);
+
+ bool tryEffectiveStall(GenericSchedulerBase::SchedCandidate &TryCand,
+ GenericSchedulerBase::SchedCandidate &Cand,
+ SchedBoundary &Zone);
+
/// Check for critical resource consumption. Prefer the candidate that uses
/// the most prioritized HardwareUnit. If both candidates use the same
/// HarwareUnit, prefer the candidate with higher priority on that
@@ -252,8 +270,6 @@ class CandidateHeuristics {
class AMDGPUCoExecSchedStrategy final : public GCNSchedStrategy {
protected:
- bool tryEffectiveStall(SchedCandidate &Cand, SchedCandidate &TryCand,
- SchedBoundary &Zone);
AMDGPU::AMDGPUSchedReason LastAMDGPUReason = AMDGPU::AMDGPUSchedReason::None;
CandidateHeuristics Heurs;
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index aa85be758559a..bf6ca9f8ad953 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -286,42 +286,6 @@ void GCNSchedStrategy::getRegisterPressures(
Pressure[AMDGPU::RegisterPressureSets::AGPR_32] = NewPressure.getAGPRNum();
}
-unsigned GCNSchedStrategy::getStructuralStallCycles(SchedBoundary &Zone,
- SUnit *SU) const {
- // Only implemented for top-down scheduling currently.
- if (!Zone.isTop() || !SU)
- return 0;
-
- MachineInstr *MI = SU->getInstr();
- unsigned CurrCycle = Zone.getCurrCycle();
- unsigned Stall = 0;
-
- // Query SchedModel for resource stalls (unbuffered resources).
- if (SchedModel->hasInstrSchedModel() && SU->hasReservedResource) {
- const MCSchedClassDesc *SC = DAG->getSchedClass(SU);
- for (const MCWriteProcResEntry &PE :
- make_range(SchedModel->getWriteProcResBegin(SC),
- SchedModel->getWriteProcResEnd(SC))) {
- unsigned NextAvail =
- Zone.getNextResourceCycle(SC, PE.ProcResourceIdx, PE.ReleaseAtCycle,
- PE.AcquireAtCycle)
- .first;
- if (NextAvail > CurrCycle)
- Stall = std::max(Stall, NextAvail - CurrCycle);
- }
- }
-
- // Query HazardRecognizer for sequence-dependent hazard penalties.
- // AMDGPUCoExecSchedStrategy installs a GCNHazardRecognizer in both
- // pre-RA (PreRA mode) and post-RA configurations.
- if (Zone.HazardRec && Zone.HazardRec->isEnabled()) {
- auto *HR = static_cast<GCNHazardRecognizer *>(Zone.HazardRec.get());
- Stall = std::max(Stall, HR->getHazardWaitStates(MI));
- }
-
- return Stall;
-}
-
void GCNSchedStrategy::initCandidate(SchedCandidate &Cand, SUnit *SU,
bool AtTop,
const RegPressureTracker &RPTracker,
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
index 29393d035cc8c..350eefb259a4a 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
@@ -60,10 +60,6 @@ class GCNSchedStrategy : public GenericScheduler {
unsigned VGPRPressure, unsigned AGPRPressure,
bool IsBottomUp);
- /// Estimate how many cycles \p SU must wait due to structural hazards at the
- /// current boundary cycle. Returns zero when no stall is required.
- unsigned getStructuralStallCycles(SchedBoundary &Zone, SUnit *SU) const;
-
/// Evaluates instructions in the pending queue using a subset of scheduling
/// heuristics.
///
diff --git a/llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll b/llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll
index 048b9abdb4c27..0dc660143f2a8 100644
--- a/llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll
+++ b/llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll
@@ -42,9 +42,6 @@ define amdgpu_kernel void @ds_wmma(ptr addrspace(3) %base, ptr addrspace(1) %out
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_mov_b32_e32 v88, s2
-; COEXEC-NEXT: s_add_co_i32 s2, s2, s1
-; COEXEC-NEXT: s_and_b32 s3, s0, exec_lo
-; COEXEC-NEXT: s_cselect_b32 s3, 1, 0
; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v88 offset:192
; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v88 offset:128
; COEXEC-NEXT: ds_load_tr16_b128 v[40:43], v88
@@ -61,6 +58,9 @@ define amdgpu_kernel void @ds_wmma(ptr addrspace(3) %base, ptr addrspace(1) %out
; COEXEC-NEXT: ds_load_tr16_b128 v[80:83], v88 offset:896
; COEXEC-NEXT: ds_load_tr16_b128 v[92:95], v88 offset:832
; COEXEC-NEXT: ds_load_tr16_b128 v[88:91], v88 offset:768
+; COEXEC-NEXT: s_add_co_i32 s2, s2, s1
+; COEXEC-NEXT: s_and_b32 s3, s0, exec_lo
+; COEXEC-NEXT: s_cselect_b32 s3, 1, 0
; COEXEC-NEXT: s_cmp_lg_u32 s3, 1
; COEXEC-NEXT: s_wait_dscnt 0xc
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[32:39], v[24:31]
@@ -76,9 +76,9 @@ define amdgpu_kernel void @ds_wmma(ptr addrspace(3) %base, ptr addrspace(1) %out
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[88:95], v[80:87], v[0:7]
; COEXEC-NEXT: s_cbranch_scc1 .LBB0_1
; COEXEC-NEXT: ; %bb.2: ; %end
+; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_mov_b32_e32 v32, 0
-; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
; COEXEC-NEXT: s_wait_kmcnt 0x0
; COEXEC-NEXT: s_clause 0x7
; COEXEC-NEXT: global_store_b128 v32, v[28:31], s[0:1] offset:16
@@ -258,9 +258,10 @@ define amdgpu_kernel void @ds_wmma_permute(ptr addrspace(3) %base, ptr addrspace
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
; COEXEC-NEXT: s_mov_b32 s6, 0
+; COEXEC-NEXT: s_clause 0x1
; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x0 nv
-; COEXEC-NEXT: v_mov_b32_e32 v0, 0
; COEXEC-NEXT: s_load_b64 s[2:3], s[4:5], 0x10 nv
+; COEXEC-NEXT: v_mov_b32_e32 v0, 0
; COEXEC-NEXT: v_dual_mov_b32 v1, v0 :: v_dual_mov_b32 v2, v0
; COEXEC-NEXT: v_dual_mov_b32 v3, v0 :: v_dual_mov_b32 v4, v0
; COEXEC-NEXT: v_dual_mov_b32 v5, v0 :: v_dual_mov_b32 v6, v0
@@ -354,8 +355,8 @@ define amdgpu_kernel void @ds_wmma_permute(ptr addrspace(3) %base, ptr addrspace
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[120:127], v[152:159], v[0:7]
; COEXEC-NEXT: s_cbranch_scc1 .LBB1_1
; COEXEC-NEXT: ; %bb.2: ; %end
-; COEXEC-NEXT: v_mov_b32_e32 v32, 0
; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
+; COEXEC-NEXT: v_mov_b32_e32 v32, 0
; COEXEC-NEXT: s_wait_kmcnt 0x0
; COEXEC-NEXT: s_clause 0x7
; COEXEC-NEXT: global_store_b128 v32, v[28:31], s[0:1] offset:16
@@ -604,4 +605,895 @@ end:
}
+define amdgpu_kernel void @ds_wmma_block_carried(ptr addrspace(3) %base, ptr addrspace(1) %out, i1 %br0, i32 %delta) local_unnamed_addr #0 {
+; COEXEC-LABEL: ds_wmma_block_carried:
+; COEXEC: ; %bb.0: ; %entry
+; COEXEC-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; COEXEC-NEXT: s_mov_b64 s[64:65], 0
+; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; COEXEC-NEXT: s_mov_b32 s8, 0
+; COEXEC-NEXT: s_load_b32 s3, s[4:5], 0x0 nv
+; COEXEC-NEXT: s_mov_b32 s9, s8
+; COEXEC-NEXT: s_mov_b32 s10, s8
+; COEXEC-NEXT: s_mov_b32 s11, s8
+; COEXEC-NEXT: s_mov_b32 s12, s8
+; COEXEC-NEXT: s_mov_b32 s13, s8
+; COEXEC-NEXT: s_mov_b32 s14, s8
+; COEXEC-NEXT: s_mov_b32 s15, s8
+; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x10 nv
+; COEXEC-NEXT: v_mov_b32_e32 v0, 0
+; COEXEC-NEXT: v_dual_mov_b32 v1, v0 :: v_dual_mov_b32 v2, v0
+; COEXEC-NEXT: v_dual_mov_b32 v3, v0 :: v_dual_mov_b32 v4, v0
+; COEXEC-NEXT: v_dual_mov_b32 v5, v0 :: v_dual_mov_b32 v6, v0
+; COEXEC-NEXT: v_dual_mov_b32 v7, v0 :: v_dual_mov_b32 v8, v0
+; COEXEC-NEXT: v_dual_mov_b32 v9, v0 :: v_dual_mov_b32 v10, v0
+; COEXEC-NEXT: v_dual_mov_b32 v11, v0 :: v_dual_mov_b32 v12, v0
+; COEXEC-NEXT: v_dual_mov_b32 v13, v0 :: v_dual_mov_b32 v14, v0
+; COEXEC-NEXT: v_dual_mov_b32 v15, v0 :: v_dual_mov_b32 v16, v0
+; COEXEC-NEXT: v_dual_mov_b32 v17, v0 :: v_dual_mov_b32 v18, v0
+; COEXEC-NEXT: v_dual_mov_b32 v19, v0 :: v_dual_mov_b32 v20, v0
+; COEXEC-NEXT: v_dual_mov_b32 v21, v0 :: v_dual_mov_b32 v22, v0
+; COEXEC-NEXT: v_dual_mov_b32 v23, v0 :: v_dual_mov_b32 v24, v0
+; COEXEC-NEXT: v_dual_mov_b32 v25, v0 :: v_dual_mov_b32 v26, v0
+; COEXEC-NEXT: v_dual_mov_b32 v27, v0 :: v_dual_mov_b32 v28, v0
+; COEXEC-NEXT: s_wait_kmcnt 0x0
+; COEXEC-NEXT: s_bitcmp1_b32 s0, 0
+; COEXEC-NEXT: v_mov_b32_e32 v29, v0
+; COEXEC-NEXT: s_cselect_b32 s2, -1, 0
+; COEXEC-NEXT: v_mov_b32_e32 v30, v0
+; COEXEC-NEXT: s_add_co_i32 s0, s3, s1
+; COEXEC-NEXT: v_mov_b32_e32 v31, v0
+; COEXEC-NEXT: s_xor_b32 s2, s2, -1
+; COEXEC-NEXT: v_mov_b32_e32 v52, s3
+; COEXEC-NEXT: ds_load_tr16_b128 v[40:43], v52
+; COEXEC-NEXT: ds_load_tr16_b128 v[44:47], v52 offset:64
+; COEXEC-NEXT: ds_load_tr16_b128 v[56:59], v52 offset:128
+; COEXEC-NEXT: ds_load_tr16_b128 v[60:63], v52 offset:192
+; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v52 offset:256
+; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v52 offset:320
+; COEXEC-NEXT: ds_load_tr16_b128 v[48:51], v52 offset:384
+; COEXEC-NEXT: ds_load_tr16_b128 v[52:55], v52 offset:448
+; COEXEC-NEXT: s_wait_dscnt 0x0
+; COEXEC-NEXT: .LBB2_1: ; %loop
+; COEXEC-NEXT: ; =>This Inner Loop Header: Depth=1
+; COEXEC-NEXT: s_wait_dscnt 0x4
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[56:63], v[24:31]
+; COEXEC-NEXT: s_wait_dscnt 0x0
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[32:39], v[48:55], v[16:23]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[56:63], v[24:31]
+; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: v_mov_b64_e32 v[70:71], s[14:15]
+; COEXEC-NEXT: v_mov_b64_e32 v[68:69], s[12:13]
+; COEXEC-NEXT: v_mov_b64_e32 v[66:67], s[10:11]
+; COEXEC-NEXT: v_mov_b64_e32 v[64:65], s[8:9]
+; COEXEC-NEXT: v_mov_b32_e32 v72, s0
+; COEXEC-NEXT: s_add_co_i32 s0, s0, s1
+; COEXEC-NEXT: s_and_b32 s3, s2, exec_lo
+; COEXEC-NEXT: s_cselect_b32 s3, 1, 0
+; COEXEC-NEXT: ds_load_tr16_b128 v[40:43], v72
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[64:71], v[8:15]
+; COEXEC-NEXT: ds_load_tr16_b128 v[44:47], v72 offset:64
+; COEXEC-NEXT: ds_load_tr16_b128 v[56:59], v72 offset:128
+; COEXEC-NEXT: ds_load_tr16_b128 v[60:63], v72 offset:192
+; COEXEC-NEXT: s_cmp_lg_u32 s3, 1
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[64:71], v[64:71], v[0:7]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[32:39], v[48:55], v[16:23]
+; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v72 offset:256
+; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v72 offset:320
+; COEXEC-NEXT: ds_load_tr16_b128 v[48:51], v72 offset:384
+; COEXEC-NEXT: ds_load_tr16_b128 v[52:55], v72 offset:448
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[64:71], v[8:15]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[64:71], v[64:71], v[0:7]
+; COEXEC-NEXT: s_cbranch_scc1 .LBB2_1
+; COEXEC-NEXT: ; %bb.2: ; %end
+; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
+; COEXEC-NEXT: s_wait_dscnt 0x3
+; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: v_mov_b32_e32 v32, 0
+; COEXEC-NEXT: s_wait_kmcnt 0x0
+; COEXEC-NEXT: s_clause 0x7
+; COEXEC-NEXT: global_store_b128 v32, v[28:31], s[0:1] offset:16
+; COEXEC-NEXT: global_store_b128 v32, v[24:27], s[0:1]
+; COEXEC-NEXT: global_store_b128 v32, v[20:23], s[0:1] offset:144
+; COEXEC-NEXT: global_store_b128 v32, v[16:19], s[0:1] offset:128
+; COEXEC-NEXT: global_store_b128 v32, v[12:15], s[0:1] offset:272
+; COEXEC-NEXT: global_store_b128 v32, v[8:11], s[0:1] offset:256
+; COEXEC-NEXT: global_store_b128 v32, v[4:7], s[0:1] offset:400
+; COEXEC-NEXT: global_store_b128 v32, v[0:3], s[0:1] offset:384
+; COEXEC-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; COEXEC-NEXT: s_endpgm
+;
+; GCN-LABEL: ds_wmma_block_carried:
+; GCN: ; %bb.0: ; %entry
+; GCN-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GCN-NEXT: s_mov_b64 s[64:65], 0
+; GCN-NEXT: v_nop
+; GCN-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GCN-NEXT: s_clause 0x1
+; GCN-NEXT: s_load_b64 s[0:1], s[4:5], 0x10 nv
+; GCN-NEXT: s_load_b32 s2, s[4:5], 0x0 nv
+; GCN-NEXT: v_mov_b32_e32 v0, 0
+; GCN-NEXT: s_mov_b32 s8, 0
+; GCN-NEXT: s_mov_b32 s9, s8
+; GCN-NEXT: s_mov_b32 s10, s8
+; GCN-NEXT: s_mov_b32 s11, s8
+; GCN-NEXT: s_mov_b32 s12, s8
+; GCN-NEXT: s_mov_b32 s13, s8
+; GCN-NEXT: s_mov_b32 s14, s8
+; GCN-NEXT: s_mov_b32 s15, s8
+; GCN-NEXT: s_wait_kmcnt 0x0
+; GCN-NEXT: s_bitcmp1_b32 s0, 0
+; GCN-NEXT: v_dual_mov_b32 v7, s2 :: v_dual_mov_b32 v1, v0
+; GCN-NEXT: v_dual_mov_b32 v2, v0 :: v_dual_mov_b32 v3, v0
+; GCN-NEXT: v_mov_b32_e32 v4, v0
+; GCN-NEXT: ds_load_tr16_b128 v[32:35], v7
+; GCN-NEXT: ds_load_tr16_b128 v[36:39], v7 offset:64
+; GCN-NEXT: ds_load_tr16_b128 v[40:43], v7 offset:128
+; GCN-NEXT: ds_load_tr16_b128 v[44:47], v7 offset:192
+; GCN-NEXT: ds_load_tr16_b128 v[48:51], v7 offset:256
+; GCN-NEXT: ds_load_tr16_b128 v[52:55], v7 offset:320
+; GCN-NEXT: ds_load_tr16_b128 v[56:59], v7 offset:384
+; GCN-NEXT: ds_load_tr16_b128 v[60:63], v7 offset:448
+; GCN-NEXT: v_dual_mov_b32 v5, v0 :: v_dual_mov_b32 v6, v0
+; GCN-NEXT: v_dual_mov_b32 v7, v0 :: v_dual_mov_b32 v8, v0
+; GCN-NEXT: v_dual_mov_b32 v9, v0 :: v_dual_mov_b32 v10, v0
+; GCN-NEXT: v_dual_mov_b32 v11, v0 :: v_dual_mov_b32 v12, v0
+; GCN-NEXT: v_dual_mov_b32 v13, v0 :: v_dual_mov_b32 v14, v0
+; GCN-NEXT: v_dual_mov_b32 v15, v0 :: v_dual_mov_b32 v16, v0
+; GCN-NEXT: v_dual_mov_b32 v17, v0 :: v_dual_mov_b32 v18, v0
+; GCN-NEXT: v_dual_mov_b32 v19, v0 :: v_dual_mov_b32 v20, v0
+; GCN-NEXT: v_dual_mov_b32 v21, v0 :: v_dual_mov_b32 v22, v0
+; GCN-NEXT: v_dual_mov_b32 v23, v0 :: v_dual_mov_b32 v24, v0
+; GCN-NEXT: v_dual_mov_b32 v25, v0 :: v_dual_mov_b32 v26, v0
+; GCN-NEXT: v_dual_mov_b32 v27, v0 :: v_dual_mov_b32 v28, v0
+; GCN-NEXT: v_dual_mov_b32 v29, v0 :: v_dual_mov_b32 v30, v0
+; GCN-NEXT: v_mov_b32_e32 v31, v0
+; GCN-NEXT: s_cselect_b32 s3, -1, 0
+; GCN-NEXT: s_add_co_i32 s0, s2, s1
+; GCN-NEXT: s_xor_b32 s2, s3, -1
+; GCN-NEXT: s_wait_dscnt 0x0
+; GCN-NEXT: .LBB2_1: ; %loop
+; GCN-NEXT: ; =>This Inner Loop Header: Depth=1
+; GCN-NEXT: s_wait_dscnt 0x0
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[56:63], v[16:23]
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_mov_b64_e32 v[70:71], s[14:15]
+; GCN-NEXT: v_mov_b64_e32 v[68:69], s[12:13]
+; GCN-NEXT: v_mov_b64_e32 v[66:67], s[10:11]
+; GCN-NEXT: v_mov_b64_e32 v[64:65], s[8:9]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[40:47], v[24:31]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[56:63], v[16:23]
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_mov_b32_e32 v60, s0
+; GCN-NEXT: s_add_co_i32 s0, s0, s1
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[40:47], v[24:31]
+; GCN-NEXT: ds_load_tr16_b128 v[32:35], v60
+; GCN-NEXT: ds_load_tr16_b128 v[36:39], v60 offset:64
+; GCN-NEXT: ds_load_tr16_b128 v[40:43], v60 offset:128
+; GCN-NEXT: ds_load_tr16_b128 v[44:47], v60 offset:192
+; GCN-NEXT: ds_load_tr16_b128 v[48:51], v60 offset:256
+; GCN-NEXT: ds_load_tr16_b128 v[52:55], v60 offset:320
+; GCN-NEXT: ds_load_tr16_b128 v[56:59], v60 offset:384
+; GCN-NEXT: ds_load_tr16_b128 v[60:63], v60 offset:448
+; GCN-NEXT: s_and_b32 s3, s2, exec_lo
+; GCN-NEXT: s_cselect_b32 s3, 1, 0
+; GCN-NEXT: s_cmp_lg_u32 s3, 1
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[64:71], v[8:15]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[64:71], v[64:71], v[0:7]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[64:71], v[8:15]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[64:71], v[64:71], v[0:7]
+; GCN-NEXT: s_cbranch_scc1 .LBB2_1
+; GCN-NEXT: ; %bb.2: ; %end
+; GCN-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
+; GCN-NEXT: s_wait_dscnt 0x7
+; GCN-NEXT: v_mov_b32_e32 v32, 0
+; GCN-NEXT: s_wait_kmcnt 0x0
+; GCN-NEXT: s_clause 0x7
+; GCN-NEXT: global_store_b128 v32, v[28:31], s[0:1] offset:16
+; GCN-NEXT: global_store_b128 v32, v[24:27], s[0:1]
+; GCN-NEXT: global_store_b128 v32, v[20:23], s[0:1] offset:144
+; GCN-NEXT: global_store_b128 v32, v[16:19], s[0:1] offset:128
+; GCN-NEXT: global_store_b128 v32, v[12:15], s[0:1] offset:272
+; GCN-NEXT: global_store_b128 v32, v[8:11], s[0:1] offset:256
+; GCN-NEXT: global_store_b128 v32, v[4:7], s[0:1] offset:400
+; GCN-NEXT: global_store_b128 v32, v[0:3], s[0:1] offset:384
+; GCN-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GCN-NEXT: s_endpgm
+entry:
+
+ %p0 = getelementptr inbounds nuw i8, ptr addrspace(3) %base, i32 0
+ %p1 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 64
+ %p2 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 128
+ %p3 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 192
+ %p4 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 256
+ %p5 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 320
+ %p6 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 384
+ %p7 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 448
+ %p8 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 512
+ %p9 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 576
+ %p10 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 640
+ %p11 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 704
+ %p12 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 768
+ %p13 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 832
+ %p14 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 896
+ %p15 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 960
+ %l0 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) %p0)
+ %l1 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p1)
+ %l2 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p2)
+ %l3 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p3)
+ %l4 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p4)
+ %l5 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p5)
+ %l6 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p6)
+ %l7 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p7)
+ %l8 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p8)
+ %l9 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p9)
+ %l10 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p10)
+ %l11 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p11)
+ %l12 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p12)
+ %l13 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p13)
+ %l14 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p14)
+ %l15 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p15)
+ %vec0 = shufflevector <8 x half> %l0, <8 x half> %l1, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec1 = shufflevector <8 x half> %l2, <8 x half> %l3, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec2 = shufflevector <8 x half> %l4, <8 x half> %l5, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec3 = shufflevector <8 x half> %l6, <8 x half> %l7, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec4 = shufflevector <8 x half> %l8, <8 x half> %l9, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec5 = shufflevector <8 x half> %l10, <8 x half> %l11, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec6 = shufflevector <8 x half> %l12, <8 x half> %l13, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec7 = shufflevector <8 x half> %l14, <8 x half> %l15, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+
+ br label %loop
+
+loop:
+ %baseOff = phi i32 [ 0, %entry ], [ %newBaseOff, %loop ]
+ %wvec0 = phi <8 x float> [ <float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0>, %entry ], [ %wmma01, %loop ]
+ %wvec1 = phi <8 x float> [ <float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0>, %entry ], [ %wmma11, %loop ]
+ %wvec2 = phi <8 x float> [ <float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0>, %entry ], [ %wmma21, %loop ]
+ %wvec3 = phi <8 x float> [ <float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0>, %entry ], [ %wmma31, %loop ]
+ %invec0 = phi <16 x half> [ %vec0, %entry ], [ %vec0l, %loop ]
+ %invec1 = phi <16 x half> [ %vec1, %entry ], [ %vec1l, %loop ]
+ %invec2 = phi <16 x half> [ %vec2, %entry ], [ %vec2l, %loop ]
+ %invec3 = phi <16 x half> [ %vec3, %entry ], [ %vec3l, %loop ]
+ %wmma00 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %invec0, i1 false, <16 x half> %invec1, i16 0, <8 x float> %wvec0, i1 false, i1 false)
+ %wmma01 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %invec0, i1 false, <16 x half> %invec1, i16 0, <8 x float> %wmma00, i1 false, i1 false)
+ %wmma10 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %invec2, i1 false, <16 x half> %invec3, i16 0, <8 x float> %wvec1, i1 false, i1 false)
+ %wmma11 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %invec2, i1 false, <16 x half> %invec3, i16 0, <8 x float> %wmma10, i1 false, i1 false)
+ %wmma20 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> zeroinitializer, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %wvec2, i1 false, i1 false)
+ %wmma21 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> zeroinitializer, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %wmma20, i1 false, i1 false)
+ %wmma30 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> zeroinitializer, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %wvec3, i1 false, i1 false)
+ %wmma31 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> zeroinitializer, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %wmma30, i1 false, i1 false)
+ %newBaseOff = or disjoint i32 %baseOff, %delta
+ %p0l = getelementptr inbounds nuw i8, ptr addrspace(3) %base, i32 %newBaseOff
+ %p1l = getelementptr inbounds nuw i8, ptr addrspace(3) %p0l, i32 64
+ %p2l = getelementptr inbounds nuw i8, ptr addrspace(3) %p0l, i32 128
+ %p3l = getelementptr inbounds nuw i8, ptr addrspace(3) %p0l, i32 192
+ %p4l = getelementptr inbounds nuw i8, ptr addrspace(3) %p0l, i32 256
+ %p5l = getelementptr inbounds nuw i8, ptr addrspace(3) %p0l, i32 320
+ %p6l = getelementptr inbounds nuw i8, ptr addrspace(3) %p0l, i32 384
+ %p7l = getelementptr inbounds nuw i8, ptr addrspace(3) %p0l, i32 448
+ %l0l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) %p0l)
+ %l1l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p1l)
+ %l2l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p2l)
+ %l3l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p3l)
+ %l4l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p4l)
+ %l5l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p5l)
+ %l6l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p6l)
+ %l7l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p7l)
+ %vec0l = shufflevector <8 x half> %l0l, <8 x half> %l1l, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec1l = shufflevector <8 x half> %l2l, <8 x half> %l3l, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec2l = shufflevector <8 x half> %l4l, <8 x half> %l5l, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec3l = shufflevector <8 x half> %l6l, <8 x half> %l7l, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+
+
+ br i1 %br0, label %loop, label %end
+
+end:
+ %out1 = getelementptr inbounds nuw i8, ptr addrspace(1) %out, i32 128
+ %out2 = getelementptr inbounds nuw i8, ptr addrspace(1) %out, i32 256
+ %out3 = getelementptr inbounds nuw i8, ptr addrspace(1) %out, i32 384
+ store <8 x float> %wmma01, ptr addrspace(1) %out, align 16
+ store <8 x float> %wmma11, ptr addrspace(1) %out1, align 16
+ store <8 x float> %wmma21, ptr addrspace(1) %out2, align 16
+ store <8 x float> %wmma31, ptr addrspace(1) %out3, align 16
+ ret void
+}
+
+define amdgpu_kernel void @ds_wmma_loop_carried(ptr addrspace(3) %base, ptr addrspace(1) %out, i1 %br0, i32 %delta) local_unnamed_addr #0 {
+; COEXEC-LABEL: ds_wmma_loop_carried:
+; COEXEC: ; %bb.0: ; %entry
+; COEXEC-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; COEXEC-NEXT: s_mov_b64 s[64:65], 0
+; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; COEXEC-NEXT: s_mov_b32 s8, 0
+; COEXEC-NEXT: s_clause 0x1
+; COEXEC-NEXT: s_load_b32 s2, s[4:5], 0x0 nv
+; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x10 nv
+; COEXEC-NEXT: s_mov_b32 s9, s8
+; COEXEC-NEXT: s_mov_b32 s10, s8
+; COEXEC-NEXT: s_mov_b32 s11, s8
+; COEXEC-NEXT: s_mov_b32 s12, s8
+; COEXEC-NEXT: s_mov_b32 s13, s8
+; COEXEC-NEXT: s_mov_b32 s14, s8
+; COEXEC-NEXT: s_mov_b32 s15, s8
+; COEXEC-NEXT: v_mov_b64_e32 v[70:71], s[14:15]
+; COEXEC-NEXT: v_mov_b64_e32 v[68:69], s[12:13]
+; COEXEC-NEXT: v_mov_b64_e32 v[66:67], s[10:11]
+; COEXEC-NEXT: v_mov_b64_e32 v[64:65], s[8:9]
+; COEXEC-NEXT: s_wait_kmcnt 0x0
+; COEXEC-NEXT: v_mov_b32_e32 v40, s2
+; COEXEC-NEXT: s_bitcmp1_b32 s0, 0
+; COEXEC-NEXT: s_cselect_b32 s3, -1, 0
+; COEXEC-NEXT: s_add_co_i32 s0, s2, s1
+; COEXEC-NEXT: ds_load_tr16_b128 v[8:11], v40
+; COEXEC-NEXT: ds_load_tr16_b128 v[12:15], v40 offset:64
+; COEXEC-NEXT: ds_load_tr16_b128 v[4:7], v40 offset:320
+; COEXEC-NEXT: ds_load_tr16_b128 v[0:3], v40 offset:256
+; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v40 offset:576
+; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v40 offset:512
+; COEXEC-NEXT: ds_load_tr16_b128 v[76:79], v40 offset:832
+; COEXEC-NEXT: ds_load_tr16_b128 v[72:75], v40 offset:768
+; COEXEC-NEXT: ds_load_tr16_b128 v[24:27], v40 offset:128
+; COEXEC-NEXT: ds_load_tr16_b128 v[28:31], v40 offset:192
+; COEXEC-NEXT: ds_load_tr16_b128 v[16:19], v40 offset:384
+; COEXEC-NEXT: ds_load_tr16_b128 v[20:23], v40 offset:448
+; COEXEC-NEXT: ds_load_tr16_b128 v[80:83], v40 offset:640
+; COEXEC-NEXT: ds_load_tr16_b128 v[84:87], v40 offset:704
+; COEXEC-NEXT: ds_load_tr16_b128 v[88:91], v40 offset:896
+; COEXEC-NEXT: ds_load_tr16_b128 v[92:95], v40 offset:960
+; COEXEC-NEXT: s_xor_b32 s2, s3, -1
+; COEXEC-NEXT: s_wait_dscnt 0xe
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[8:15], v[64:71], 0
+; COEXEC-NEXT: s_wait_dscnt 0xc
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[48:55], v[0:7], v[64:71], 0
+; COEXEC-NEXT: s_wait_dscnt 0xa
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[40:47], v[32:39], v[64:71], 0
+; COEXEC-NEXT: s_wait_dscnt 0x8
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[72:79], v[64:71], 0
+; COEXEC-NEXT: s_wait_dscnt 0x6
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[24:31], v[64:71], v[56:63]
+; COEXEC-NEXT: s_wait_dscnt 0x4
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[48:55], v[16:23], v[64:71], v[48:55]
+; COEXEC-NEXT: s_wait_dscnt 0x2
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[40:47], v[80:87], v[64:71], v[40:47]
+; COEXEC-NEXT: s_wait_dscnt 0x0
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[88:95], v[64:71], v[32:39]
+; COEXEC-NEXT: .LBB3_1: ; %loop
+; COEXEC-NEXT: ; =>This Inner Loop Header: Depth=1
+; COEXEC-NEXT: s_wait_dscnt 0x4
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[8:15], v[24:31], v[56:63]
+; COEXEC-NEXT: s_wait_dscnt 0x0
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[48:55], v[0:7], v[16:23], v[48:55]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[8:15], v[24:31], v[56:63]
+; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: v_mov_b64_e32 v[70:71], s[14:15]
+; COEXEC-NEXT: v_mov_b64_e32 v[68:69], s[12:13]
+; COEXEC-NEXT: v_mov_b64_e32 v[66:67], s[10:11]
+; COEXEC-NEXT: v_mov_b64_e32 v[64:65], s[8:9]
+; COEXEC-NEXT: v_mov_b32_e32 v72, s0
+; COEXEC-NEXT: s_add_co_i32 s0, s0, s1
+; COEXEC-NEXT: s_and_b32 s3, s2, exec_lo
+; COEXEC-NEXT: s_cselect_b32 s3, 1, 0
+; COEXEC-NEXT: ds_load_tr16_b128 v[8:11], v72
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[40:47], v[64:71], v[64:71], v[40:47]
+; COEXEC-NEXT: ds_load_tr16_b128 v[12:15], v72 offset:64
+; COEXEC-NEXT: ds_load_tr16_b128 v[24:27], v72 offset:128
+; COEXEC-NEXT: ds_load_tr16_b128 v[28:31], v72 offset:192
+; COEXEC-NEXT: s_cmp_lg_u32 s3, 1
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[64:71], v[64:71], v[32:39]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[48:55], v[0:7], v[16:23], v[48:55]
+; COEXEC-NEXT: ds_load_tr16_b128 v[0:3], v72 offset:256
+; COEXEC-NEXT: ds_load_tr16_b128 v[4:7], v72 offset:320
+; COEXEC-NEXT: ds_load_tr16_b128 v[16:19], v72 offset:384
+; COEXEC-NEXT: ds_load_tr16_b128 v[20:23], v72 offset:448
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[40:47], v[64:71], v[64:71], v[40:47]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[64:71], v[64:71], v[32:39]
+; COEXEC-NEXT: s_cbranch_scc1 .LBB3_1
+; COEXEC-NEXT: ; %bb.2: ; %end
+; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
+; COEXEC-NEXT: s_wait_dscnt 0x3
+; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: v_mov_b32_e32 v0, 0
+; COEXEC-NEXT: s_wait_kmcnt 0x0
+; COEXEC-NEXT: s_clause 0x7
+; COEXEC-NEXT: global_store_b128 v0, v[60:63], s[0:1] offset:16
+; COEXEC-NEXT: global_store_b128 v0, v[56:59], s[0:1]
+; COEXEC-NEXT: global_store_b128 v0, v[52:55], s[0:1] offset:144
+; COEXEC-NEXT: global_store_b128 v0, v[48:51], s[0:1] offset:128
+; COEXEC-NEXT: global_store_b128 v0, v[44:47], s[0:1] offset:272
+; COEXEC-NEXT: global_store_b128 v0, v[40:43], s[0:1] offset:256
+; COEXEC-NEXT: global_store_b128 v0, v[36:39], s[0:1] offset:400
+; COEXEC-NEXT: global_store_b128 v0, v[32:35], s[0:1] offset:384
+; COEXEC-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; COEXEC-NEXT: s_endpgm
+;
+; GCN-LABEL: ds_wmma_loop_carried:
+; GCN: ; %bb.0: ; %entry
+; GCN-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GCN-NEXT: s_mov_b64 s[64:65], 0
+; GCN-NEXT: v_nop
+; GCN-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GCN-NEXT: s_clause 0x1
+; GCN-NEXT: s_load_b32 s2, s[4:5], 0x0 nv
+; GCN-NEXT: s_load_b64 s[0:1], s[4:5], 0x10 nv
+; GCN-NEXT: s_mov_b32 s8, 0
+; GCN-NEXT: s_mov_b32 s9, s8
+; GCN-NEXT: s_mov_b32 s10, s8
+; GCN-NEXT: s_mov_b32 s11, s8
+; GCN-NEXT: s_mov_b32 s12, s8
+; GCN-NEXT: s_mov_b32 s13, s8
+; GCN-NEXT: s_mov_b32 s14, s8
+; GCN-NEXT: s_mov_b32 s15, s8
+; GCN-NEXT: v_mov_b64_e32 v[68:69], s[12:13]
+; GCN-NEXT: v_mov_b64_e32 v[70:71], s[14:15]
+; GCN-NEXT: v_mov_b64_e32 v[66:67], s[10:11]
+; GCN-NEXT: v_mov_b64_e32 v[64:65], s[8:9]
+; GCN-NEXT: s_wait_kmcnt 0x0
+; GCN-NEXT: v_mov_b32_e32 v80, s2
+; GCN-NEXT: s_bitcmp1_b32 s0, 0
+; GCN-NEXT: ds_load_tr16_b128 v[8:11], v80 offset:512
+; GCN-NEXT: ds_load_tr16_b128 v[12:15], v80 offset:576
+; GCN-NEXT: ds_load_tr16_b128 v[48:51], v80 offset:768
+; GCN-NEXT: ds_load_tr16_b128 v[52:55], v80 offset:832
+; GCN-NEXT: ds_load_tr16_b128 v[72:75], v80 offset:640
+; GCN-NEXT: ds_load_tr16_b128 v[76:79], v80 offset:704
+; GCN-NEXT: ds_load_tr16_b128 v[0:3], v80
+; GCN-NEXT: ds_load_tr16_b128 v[4:7], v80 offset:64
+; GCN-NEXT: ds_load_tr16_b128 v[40:43], v80 offset:256
+; GCN-NEXT: ds_load_tr16_b128 v[44:47], v80 offset:320
+; GCN-NEXT: ds_load_tr16_b128 v[56:59], v80 offset:384
+; GCN-NEXT: ds_load_tr16_b128 v[60:63], v80 offset:448
+; GCN-NEXT: s_cselect_b32 s3, -1, 0
+; GCN-NEXT: s_add_co_i32 s0, s2, s1
+; GCN-NEXT: s_xor_b32 s2, s3, -1
+; GCN-NEXT: s_wait_dscnt 0xa
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[8:15], v[64:71], 0
+; GCN-NEXT: s_wait_dscnt 0x8
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[48:55], v[64:71], 0
+; GCN-NEXT: ds_load_tr16_b128 v[48:51], v80 offset:128
+; GCN-NEXT: ds_load_tr16_b128 v[52:55], v80 offset:192
+; GCN-NEXT: s_wait_dscnt 0x8
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[72:79], v[64:71], v[16:23]
+; GCN-NEXT: ds_load_tr16_b128 v[72:75], v80 offset:896
+; GCN-NEXT: ds_load_tr16_b128 v[76:79], v80 offset:960
+; GCN-NEXT: s_wait_dscnt 0x8
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[0:7], v[64:71], 0
+; GCN-NEXT: s_wait_dscnt 0x6
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[64:71], 0
+; GCN-NEXT: s_wait_dscnt 0x2
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[48:55], v[64:71], v[32:39]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[56:63], v[64:71], v[24:31]
+; GCN-NEXT: s_wait_dscnt 0x0
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[72:79], v[64:71], v[8:15]
+; GCN-NEXT: .LBB3_1: ; %loop
+; GCN-NEXT: ; =>This Inner Loop Header: Depth=1
+; GCN-NEXT: s_wait_dscnt 0x0
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[56:63], v[24:31]
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_mov_b64_e32 v[70:71], s[14:15]
+; GCN-NEXT: v_mov_b64_e32 v[68:69], s[12:13]
+; GCN-NEXT: v_mov_b64_e32 v[66:67], s[10:11]
+; GCN-NEXT: v_mov_b64_e32 v[64:65], s[8:9]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[0:7], v[48:55], v[32:39]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[56:63], v[24:31]
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_mov_b32_e32 v60, s0
+; GCN-NEXT: s_add_co_i32 s0, s0, s1
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[0:7], v[48:55], v[32:39]
+; GCN-NEXT: ds_load_tr16_b128 v[0:3], v60
+; GCN-NEXT: ds_load_tr16_b128 v[4:7], v60 offset:64
+; GCN-NEXT: ds_load_tr16_b128 v[48:51], v60 offset:128
+; GCN-NEXT: ds_load_tr16_b128 v[52:55], v60 offset:192
+; GCN-NEXT: ds_load_tr16_b128 v[40:43], v60 offset:256
+; GCN-NEXT: ds_load_tr16_b128 v[44:47], v60 offset:320
+; GCN-NEXT: ds_load_tr16_b128 v[56:59], v60 offset:384
+; GCN-NEXT: ds_load_tr16_b128 v[60:63], v60 offset:448
+; GCN-NEXT: s_and_b32 s3, s2, exec_lo
+; GCN-NEXT: s_cselect_b32 s3, 1, 0
+; GCN-NEXT: s_cmp_lg_u32 s3, 1
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[64:71], v[64:71], v[16:23]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[64:71], v[8:15]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[64:71], v[64:71], v[16:23]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[64:71], v[8:15]
+; GCN-NEXT: s_cbranch_scc1 .LBB3_1
+; GCN-NEXT: ; %bb.2: ; %end
+; GCN-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
+; GCN-NEXT: s_wait_dscnt 0x7
+; GCN-NEXT: v_mov_b32_e32 v0, 0
+; GCN-NEXT: s_wait_kmcnt 0x0
+; GCN-NEXT: s_clause 0x7
+; GCN-NEXT: global_store_b128 v0, v[36:39], s[0:1] offset:16
+; GCN-NEXT: global_store_b128 v0, v[32:35], s[0:1]
+; GCN-NEXT: global_store_b128 v0, v[28:31], s[0:1] offset:144
+; GCN-NEXT: global_store_b128 v0, v[24:27], s[0:1] offset:128
+; GCN-NEXT: global_store_b128 v0, v[20:23], s[0:1] offset:272
+; GCN-NEXT: global_store_b128 v0, v[16:19], s[0:1] offset:256
+; GCN-NEXT: global_store_b128 v0, v[12:15], s[0:1] offset:400
+; GCN-NEXT: global_store_b128 v0, v[8:11], s[0:1] offset:384
+; GCN-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GCN-NEXT: s_endpgm
+entry:
+
+ %p0 = getelementptr inbounds nuw i8, ptr addrspace(3) %base, i32 0
+ %p1 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 64
+ %p2 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 128
+ %p3 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 192
+ %p4 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 256
+ %p5 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 320
+ %p6 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 384
+ %p7 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 448
+ %p8 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 512
+ %p9 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 576
+ %p10 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 640
+ %p11 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 704
+ %p12 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 768
+ %p13 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 832
+ %p14 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 896
+ %p15 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 960
+ %l0 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) %p0)
+ %l1 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p1)
+ %l2 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p2)
+ %l3 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p3)
+ %l4 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p4)
+ %l5 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p5)
+ %l6 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p6)
+ %l7 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p7)
+ %l8 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p8)
+ %l9 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p9)
+ %l10 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p10)
+ %l11 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p11)
+ %l12 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p12)
+ %l13 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p13)
+ %l14 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p14)
+ %l15 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p15)
+ %vec0 = shufflevector <8 x half> %l0, <8 x half> %l1, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec1 = shufflevector <8 x half> %l2, <8 x half> %l3, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec2 = shufflevector <8 x half> %l4, <8 x half> %l5, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec3 = shufflevector <8 x half> %l6, <8 x half> %l7, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec4 = shufflevector <8 x half> %l8, <8 x half> %l9, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec5 = shufflevector <8 x half> %l10, <8 x half> %l11, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec6 = shufflevector <8 x half> %l12, <8 x half> %l13, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec7 = shufflevector <8 x half> %l14, <8 x half> %l15, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %incwmma00 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec0, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> zeroinitializer, i1 false, i1 false)
+ %incwmma01 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec1, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %incwmma00, i1 false, i1 false)
+ %incwmma10 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec2, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> zeroinitializer, i1 false, i1 false)
+ %incwmma11 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec3, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %incwmma10, i1 false, i1 false)
+ %incwmma20 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec4, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> zeroinitializer, i1 false, i1 false)
+ %incwmma21 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec5, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %incwmma20, i1 false, i1 false)
+ %incwmma30 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec6, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> zeroinitializer, i1 false, i1 false)
+ %incwmma31 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec7, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %incwmma30, i1 false, i1 false)
+
+
+
+ br label %loop
+
+loop:
+ %baseOff = phi i32 [ 0, %entry ], [ %newBaseOff, %loop ]
+ %wvec0 = phi <8 x float> [ %incwmma01, %entry ], [ %wmma01, %loop ]
+ %wvec1 = phi <8 x float> [ %incwmma11, %entry ], [ %wmma11, %loop ]
+ %wvec2 = phi <8 x float> [ %incwmma21, %entry ], [ %wmma21, %loop ]
+ %wvec3 = phi <8 x float> [ %incwmma31, %entry ], [ %wmma31, %loop ]
+ %invec0 = phi <16 x half> [ %vec0, %entry ], [ %vec0l, %loop ]
+ %invec1 = phi <16 x half> [ %vec1, %entry ], [ %vec1l, %loop ]
+ %invec2 = phi <16 x half> [ %vec2, %entry ], [ %vec2l, %loop ]
+ %invec3 = phi <16 x half> [ %vec3, %entry ], [ %vec3l, %loop ]
+ %wmma00 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %invec0, i1 false, <16 x half> %invec1, i16 0, <8 x float> %wvec0, i1 false, i1 false)
+ %wmma01 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %invec0, i1 false, <16 x half> %invec1, i16 0, <8 x float> %wmma00, i1 false, i1 false)
+ %wmma10 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %invec2, i1 false, <16 x half> %invec3, i16 0, <8 x float> %wvec1, i1 false, i1 false)
+ %wmma11 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %invec2, i1 false, <16 x half> %invec3, i16 0, <8 x float> %wmma10, i1 false, i1 false)
+ %wmma20 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> zeroinitializer, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %wvec2, i1 false, i1 false)
+ %wmma21 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> zeroinitializer, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %wmma20, i1 false, i1 false)
+ %wmma30 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> zeroinitializer, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %wvec3, i1 false, i1 false)
+ %wmma31 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> zeroinitializer, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %wmma30, i1 false, i1 false)
+ %newBaseOff = or disjoint i32 %baseOff, %delta
+ %p0l = getelementptr inbounds nuw i8, ptr addrspace(3) %base, i32 %newBaseOff
+ %p1l = getelementptr inbounds nuw i8, ptr addrspace(3) %p0l, i32 64
+ %p2l = getelementptr inbounds nuw i8, ptr addrspace(3) %p0l, i32 128
+ %p3l = getelementptr inbounds nuw i8, ptr addrspace(3) %p0l, i32 192
+ %p4l = getelementptr inbounds nuw i8, ptr addrspace(3) %p0l, i32 256
+ %p5l = getelementptr inbounds nuw i8, ptr addrspace(3) %p0l, i32 320
+ %p6l = getelementptr inbounds nuw i8, ptr addrspace(3) %p0l, i32 384
+ %p7l = getelementptr inbounds nuw i8, ptr addrspace(3) %p0l, i32 448
+ %l0l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) %p0l)
+ %l1l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p1l)
+ %l2l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p2l)
+ %l3l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p3l)
+ %l4l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p4l)
+ %l5l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p5l)
+ %l6l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p6l)
+ %l7l = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p7l)
+ %vec0l = shufflevector <8 x half> %l0l, <8 x half> %l1l, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec1l = shufflevector <8 x half> %l2l, <8 x half> %l3l, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec2l = shufflevector <8 x half> %l4l, <8 x half> %l5l, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec3l = shufflevector <8 x half> %l6l, <8 x half> %l7l, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+
+
+ br i1 %br0, label %loop, label %end
+
+end:
+ %out1 = getelementptr inbounds nuw i8, ptr addrspace(1) %out, i32 128
+ %out2 = getelementptr inbounds nuw i8, ptr addrspace(1) %out, i32 256
+ %out3 = getelementptr inbounds nuw i8, ptr addrspace(1) %out, i32 384
+ store <8 x float> %wmma01, ptr addrspace(1) %out, align 16
+ store <8 x float> %wmma11, ptr addrspace(1) %out1, align 16
+ store <8 x float> %wmma21, ptr addrspace(1) %out2, align 16
+ store <8 x float> %wmma31, ptr addrspace(1) %out3, align 16
+ ret void
+}
+
+define amdgpu_kernel void @ds_wmma_no_block_carried(ptr addrspace(3) %base, ptr addrspace(1) %out, i1 %br0, i32 %delta) local_unnamed_addr #0 {
+; COEXEC-LABEL: ds_wmma_no_block_carried:
+; COEXEC: ; %bb.0: ; %entry
+; COEXEC-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; COEXEC-NEXT: s_mov_b64 s[64:65], 0
+; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; COEXEC-NEXT: s_mov_b32 s8, 0
+; COEXEC-NEXT: s_clause 0x1
+; COEXEC-NEXT: s_load_b32 s0, s[4:5], 0x0 nv
+; COEXEC-NEXT: s_load_b32 s1, s[4:5], 0x10 nv
+; COEXEC-NEXT: s_mov_b32 s9, s8
+; COEXEC-NEXT: s_mov_b32 s10, s8
+; COEXEC-NEXT: s_mov_b32 s11, s8
+; COEXEC-NEXT: s_mov_b32 s12, s8
+; COEXEC-NEXT: s_mov_b32 s13, s8
+; COEXEC-NEXT: s_mov_b32 s14, s8
+; COEXEC-NEXT: s_mov_b32 s15, s8
+; COEXEC-NEXT: v_mov_b64_e32 v[102:103], s[14:15]
+; COEXEC-NEXT: v_mov_b64_e32 v[100:101], s[12:13]
+; COEXEC-NEXT: v_mov_b64_e32 v[98:99], s[10:11]
+; COEXEC-NEXT: v_mov_b64_e32 v[96:97], s[8:9]
+; COEXEC-NEXT: s_wait_kmcnt 0x0
+; COEXEC-NEXT: v_mov_b32_e32 v60, s0
+; COEXEC-NEXT: s_bitcmp1_b32 s1, 0
+; COEXEC-NEXT: s_cselect_b32 s0, -1, 0
+; COEXEC-NEXT: ds_load_tr16_b128 v[0:3], v60
+; COEXEC-NEXT: ds_load_tr16_b128 v[4:7], v60 offset:64
+; COEXEC-NEXT: ds_load_tr16_b128 v[12:15], v60 offset:320
+; COEXEC-NEXT: ds_load_tr16_b128 v[8:11], v60 offset:256
+; COEXEC-NEXT: ds_load_tr16_b128 v[20:23], v60 offset:576
+; COEXEC-NEXT: ds_load_tr16_b128 v[16:19], v60 offset:512
+; COEXEC-NEXT: ds_load_tr16_b128 v[28:31], v60 offset:832
+; COEXEC-NEXT: ds_load_tr16_b128 v[24:27], v60 offset:768
+; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v60 offset:128
+; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v60 offset:192
+; COEXEC-NEXT: ds_load_tr16_b128 v[40:43], v60 offset:384
+; COEXEC-NEXT: ds_load_tr16_b128 v[44:47], v60 offset:448
+; COEXEC-NEXT: ds_load_tr16_b128 v[48:51], v60 offset:640
+; COEXEC-NEXT: ds_load_tr16_b128 v[52:55], v60 offset:704
+; COEXEC-NEXT: ds_load_tr16_b128 v[56:59], v60 offset:896
+; COEXEC-NEXT: ds_load_tr16_b128 v[60:63], v60 offset:960
+; COEXEC-NEXT: s_xor_b32 s0, s0, -1
+; COEXEC-NEXT: s_wait_dscnt 0xe
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[88:95], v[0:7], v[96:103], 0
+; COEXEC-NEXT: s_wait_dscnt 0xc
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[80:87], v[8:15], v[96:103], 0
+; COEXEC-NEXT: s_wait_dscnt 0xa
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[72:79], v[16:23], v[96:103], 0
+; COEXEC-NEXT: s_wait_dscnt 0x8
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[64:71], v[24:31], v[96:103], 0
+; COEXEC-NEXT: s_wait_dscnt 0x6
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[88:95], v[32:39], v[96:103], v[88:95]
+; COEXEC-NEXT: s_wait_dscnt 0x4
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[80:87], v[40:47], v[96:103], v[80:87]
+; COEXEC-NEXT: s_wait_dscnt 0x2
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[72:79], v[48:55], v[96:103], v[72:79]
+; COEXEC-NEXT: s_wait_dscnt 0x0
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[64:71], v[56:63], v[96:103], v[64:71]
+; COEXEC-NEXT: .LBB4_1: ; %loop
+; COEXEC-NEXT: ; =>This Inner Loop Header: Depth=1
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[88:95], v[0:7], v[32:39], v[88:95]
+; COEXEC-NEXT: s_and_b32 s1, s0, exec_lo
+; COEXEC-NEXT: s_cselect_b32 s1, 1, 0
+; COEXEC-NEXT: s_cmp_lg_u32 s1, 1
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[80:87], v[8:15], v[40:47], v[80:87]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[72:79], v[16:23], v[48:55], v[72:79]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[64:71], v[24:31], v[56:63], v[64:71]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[88:95], v[0:7], v[32:39], v[88:95]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[80:87], v[8:15], v[40:47], v[80:87]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[72:79], v[16:23], v[48:55], v[72:79]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[64:71], v[24:31], v[56:63], v[64:71]
+; COEXEC-NEXT: s_cbranch_scc1 .LBB4_1
+; COEXEC-NEXT: ; %bb.2: ; %end
+; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
+; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: v_mov_b32_e32 v0, 0
+; COEXEC-NEXT: s_wait_kmcnt 0x0
+; COEXEC-NEXT: s_clause 0x7
+; COEXEC-NEXT: global_store_b128 v0, v[92:95], s[0:1] offset:16
+; COEXEC-NEXT: global_store_b128 v0, v[88:91], s[0:1]
+; COEXEC-NEXT: global_store_b128 v0, v[84:87], s[0:1] offset:144
+; COEXEC-NEXT: global_store_b128 v0, v[80:83], s[0:1] offset:128
+; COEXEC-NEXT: global_store_b128 v0, v[76:79], s[0:1] offset:272
+; COEXEC-NEXT: global_store_b128 v0, v[72:75], s[0:1] offset:256
+; COEXEC-NEXT: global_store_b128 v0, v[68:71], s[0:1] offset:400
+; COEXEC-NEXT: global_store_b128 v0, v[64:67], s[0:1] offset:384
+; COEXEC-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; COEXEC-NEXT: s_endpgm
+;
+; GCN-LABEL: ds_wmma_no_block_carried:
+; GCN: ; %bb.0: ; %entry
+; GCN-NEXT: s_setreg_imm32_b32 hwreg(HW_REG_WAVE_MODE, 25, 1), 1 ; msbs: dst=0 src0=0 src1=0 src2=0
+; GCN-NEXT: s_mov_b64 s[64:65], 0
+; GCN-NEXT: v_nop
+; GCN-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
+; GCN-NEXT: s_clause 0x1
+; GCN-NEXT: s_load_b32 s0, s[4:5], 0x0 nv
+; GCN-NEXT: s_load_b32 s1, s[4:5], 0x10 nv
+; GCN-NEXT: s_mov_b32 s8, 0
+; GCN-NEXT: s_mov_b32 s9, s8
+; GCN-NEXT: s_mov_b32 s10, s8
+; GCN-NEXT: s_mov_b32 s11, s8
+; GCN-NEXT: s_mov_b32 s12, s8
+; GCN-NEXT: s_mov_b32 s13, s8
+; GCN-NEXT: s_mov_b32 s14, s8
+; GCN-NEXT: s_mov_b32 s15, s8
+; GCN-NEXT: v_mov_b64_e32 v[100:101], s[12:13]
+; GCN-NEXT: v_mov_b64_e32 v[102:103], s[14:15]
+; GCN-NEXT: v_mov_b64_e32 v[98:99], s[10:11]
+; GCN-NEXT: v_mov_b64_e32 v[96:97], s[8:9]
+; GCN-NEXT: s_wait_kmcnt 0x0
+; GCN-NEXT: v_mov_b32_e32 v92, s0
+; GCN-NEXT: s_bitcmp1_b32 s1, 0
+; GCN-NEXT: ds_load_tr16_b128 v[0:3], v92
+; GCN-NEXT: ds_load_tr16_b128 v[4:7], v92 offset:64
+; GCN-NEXT: ds_load_tr16_b128 v[40:43], v92 offset:256
+; GCN-NEXT: ds_load_tr16_b128 v[44:47], v92 offset:320
+; GCN-NEXT: ds_load_tr16_b128 v[48:51], v92 offset:512
+; GCN-NEXT: ds_load_tr16_b128 v[52:55], v92 offset:576
+; GCN-NEXT: ds_load_tr16_b128 v[56:59], v92 offset:768
+; GCN-NEXT: ds_load_tr16_b128 v[60:63], v92 offset:832
+; GCN-NEXT: ds_load_tr16_b128 v[64:67], v92 offset:128
+; GCN-NEXT: ds_load_tr16_b128 v[68:71], v92 offset:192
+; GCN-NEXT: ds_load_tr16_b128 v[72:75], v92 offset:384
+; GCN-NEXT: ds_load_tr16_b128 v[76:79], v92 offset:448
+; GCN-NEXT: ds_load_tr16_b128 v[80:83], v92 offset:640
+; GCN-NEXT: ds_load_tr16_b128 v[84:87], v92 offset:704
+; GCN-NEXT: ds_load_tr16_b128 v[88:91], v92 offset:896
+; GCN-NEXT: ds_load_tr16_b128 v[92:95], v92 offset:960
+; GCN-NEXT: s_cselect_b32 s0, -1, 0
+; GCN-NEXT: s_wait_dscnt 0xe
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[0:7], v[96:103], 0
+; GCN-NEXT: s_xor_b32 s0, s0, -1
+; GCN-NEXT: s_wait_dscnt 0xc
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[96:103], 0
+; GCN-NEXT: s_wait_dscnt 0xa
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[96:103], 0
+; GCN-NEXT: s_wait_dscnt 0x8
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[56:63], v[96:103], 0
+; GCN-NEXT: s_wait_dscnt 0x6
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[64:71], v[96:103], v[32:39]
+; GCN-NEXT: s_wait_dscnt 0x4
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[72:79], v[96:103], v[24:31]
+; GCN-NEXT: s_wait_dscnt 0x2
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[80:87], v[96:103], v[16:23]
+; GCN-NEXT: s_wait_dscnt 0x0
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[88:95], v[96:103], v[8:15]
+; GCN-NEXT: .LBB4_1: ; %loop
+; GCN-NEXT: ; =>This Inner Loop Header: Depth=1
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[0:7], v[64:71], v[32:39]
+; GCN-NEXT: s_and_b32 s1, s0, exec_lo
+; GCN-NEXT: s_cselect_b32 s1, 1, 0
+; GCN-NEXT: s_cmp_lg_u32 s1, 1
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[72:79], v[24:31]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[80:87], v[16:23]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[56:63], v[88:95], v[8:15]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[0:7], v[64:71], v[32:39]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[72:79], v[24:31]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[80:87], v[16:23]
+; GCN-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[56:63], v[88:95], v[8:15]
+; GCN-NEXT: s_cbranch_scc1 .LBB4_1
+; GCN-NEXT: ; %bb.2: ; %end
+; GCN-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
+; GCN-NEXT: v_nop
+; GCN-NEXT: v_mov_b32_e32 v0, 0
+; GCN-NEXT: s_wait_kmcnt 0x0
+; GCN-NEXT: s_clause 0x7
+; GCN-NEXT: global_store_b128 v0, v[36:39], s[0:1] offset:16
+; GCN-NEXT: global_store_b128 v0, v[32:35], s[0:1]
+; GCN-NEXT: global_store_b128 v0, v[28:31], s[0:1] offset:144
+; GCN-NEXT: global_store_b128 v0, v[24:27], s[0:1] offset:128
+; GCN-NEXT: global_store_b128 v0, v[20:23], s[0:1] offset:272
+; GCN-NEXT: global_store_b128 v0, v[16:19], s[0:1] offset:256
+; GCN-NEXT: global_store_b128 v0, v[12:15], s[0:1] offset:400
+; GCN-NEXT: global_store_b128 v0, v[8:11], s[0:1] offset:384
+; GCN-NEXT: s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
+; GCN-NEXT: s_endpgm
+entry:
+
+ %p0 = getelementptr inbounds nuw i8, ptr addrspace(3) %base, i32 0
+ %p1 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 64
+ %p2 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 128
+ %p3 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 192
+ %p4 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 256
+ %p5 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 320
+ %p6 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 384
+ %p7 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 448
+ %p8 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 512
+ %p9 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 576
+ %p10 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 640
+ %p11 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 704
+ %p12 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 768
+ %p13 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 832
+ %p14 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 896
+ %p15 = getelementptr inbounds nuw i8, ptr addrspace(3) %p0, i32 960
+ %l0 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) %p0)
+ %l1 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p1)
+ %l2 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p2)
+ %l3 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p3)
+ %l4 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p4)
+ %l5 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p5)
+ %l6 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p6)
+ %l7 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p7)
+ %l8 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p8)
+ %l9 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p9)
+ %l10 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p10)
+ %l11 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p11)
+ %l12 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p12)
+ %l13 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p13)
+ %l14 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p14)
+ %l15 = tail call <8 x half> @llvm.amdgcn.ds.load.tr16.b128.v8f16(ptr addrspace(3) nonnull %p15)
+ %vec0 = shufflevector <8 x half> %l0, <8 x half> %l1, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec1 = shufflevector <8 x half> %l2, <8 x half> %l3, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec2 = shufflevector <8 x half> %l4, <8 x half> %l5, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec3 = shufflevector <8 x half> %l6, <8 x half> %l7, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec4 = shufflevector <8 x half> %l8, <8 x half> %l9, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec5 = shufflevector <8 x half> %l10, <8 x half> %l11, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec6 = shufflevector <8 x half> %l12, <8 x half> %l13, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %vec7 = shufflevector <8 x half> %l14, <8 x half> %l15, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %incwmma00 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec0, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> zeroinitializer, i1 false, i1 false)
+ %incwmma01 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec1, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %incwmma00, i1 false, i1 false)
+ %incwmma10 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec2, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> zeroinitializer, i1 false, i1 false)
+ %incwmma11 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec3, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %incwmma10, i1 false, i1 false)
+ %incwmma20 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec4, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> zeroinitializer, i1 false, i1 false)
+ %incwmma21 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec5, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %incwmma20, i1 false, i1 false)
+ %incwmma30 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec6, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> zeroinitializer, i1 false, i1 false)
+ %incwmma31 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec7, i1 false, <16 x half> zeroinitializer, i16 0, <8 x float> %incwmma30, i1 false, i1 false)
+
+
+
+ br label %loop
+
+loop:
+ %baseOff = phi i32 [ 0, %entry ], [ %newBaseOff, %loop ]
+ %wvec0 = phi <8 x float> [ %incwmma01, %entry ], [ %wmma01, %loop ]
+ %wvec1 = phi <8 x float> [ %incwmma11, %entry ], [ %wmma11, %loop ]
+ %wvec2 = phi <8 x float> [ %incwmma21, %entry ], [ %wmma21, %loop ]
+ %wvec3 = phi <8 x float> [ %incwmma31, %entry ], [ %wmma31, %loop ]
+ %wmma00 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec0, i1 false, <16 x half> %vec1, i16 0, <8 x float> %wvec0, i1 false, i1 false)
+ %wmma01 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec0, i1 false, <16 x half> %vec1, i16 0, <8 x float> %wmma00, i1 false, i1 false)
+ %wmma10 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec2, i1 false, <16 x half> %vec3, i16 0, <8 x float> %wvec1, i1 false, i1 false)
+ %wmma11 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec2, i1 false, <16 x half> %vec3, i16 0, <8 x float> %wmma10, i1 false, i1 false)
+ %wmma20 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec4, i1 false, <16 x half> %vec5, i16 0, <8 x float> %wvec2, i1 false, i1 false)
+ %wmma21 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec4, i1 false, <16 x half> %vec5, i16 0, <8 x float> %wmma20, i1 false, i1 false)
+ %wmma30 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec6, i1 false, <16 x half> %vec7, i16 0, <8 x float> %wvec3, i1 false, i1 false)
+ %wmma31 = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.f16.v8f32.v16f16(i1 false, <16 x half> %vec6, i1 false, <16 x half> %vec7, i16 0, <8 x float> %wmma30, i1 false, i1 false)
+ %newBaseOff = or disjoint i32 %baseOff, %delta
+ br i1 %br0, label %loop, label %end
+
+end:
+ %out1 = getelementptr inbounds nuw i8, ptr addrspace(1) %out, i32 128
+ %out2 = getelementptr inbounds nuw i8, ptr addrspace(1) %out, i32 256
+ %out3 = getelementptr inbounds nuw i8, ptr addrspace(1) %out, i32 384
+ store <8 x float> %wmma01, ptr addrspace(1) %out, align 16
+ store <8 x float> %wmma11, ptr addrspace(1) %out1, align 16
+ store <8 x float> %wmma21, ptr addrspace(1) %out2, align 16
+ store <8 x float> %wmma31, ptr addrspace(1) %out3, align 16
+ ret void
+}
+
+
attributes #0 = { "amdgpu-flat-work-group-size"="32,32" "amdgpu-waves-per-eu"="1,1" }
>From d3fd6b354980faa8fe63c1b2df8f2a384f025832 Mon Sep 17 00:00:00 2001
From: Jeffrey Byrnes <Jeffrey.Byrnes at amd.com>
Date: Thu, 28 May 2026 11:41:03 -0700
Subject: [PATCH 02/10] Claude Code review
Change-Id: Iab06de2981b27667cc29a56931dd378ecf7a1b0c
---
.../AMDGPU/AMDGPUCoExecSchedStrategy.cpp | 50 ++--
.../Target/AMDGPU/AMDGPUCoExecSchedStrategy.h | 5 +
llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll | 214 +++++++++---------
3 files changed, 135 insertions(+), 134 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
index 9f23ddb7cf6bb..3311a4432a102 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
@@ -544,17 +544,13 @@ CandidateHeuristics::getHWUIFromFlavor(InstructionFlavor Flavor) {
return nullptr;
}
-unsigned CandidateHeuristics::getHWUICyclesForSU(SUnit *SU) {
- assert(SchedModel && SchedModel->hasInstrSchedModel());
- MachineInstr *MI = SU->getInstr();
- if (SII->isDS(*MI))
- return SchedModel->computeInstrLatency(MI);
-
+unsigned CandidateHeuristics::getMaxBlockingCycles(const MCSchedClassDesc *SC,
+ const MachineInstr *MI) {
// Loads and stores are not pipelined
if (MI->mayLoadOrStore())
return SchedModel->computeInstrLatency(MI, false);
+
unsigned ReleaseAtCycle = 0;
- const MCSchedClassDesc *SC = DAG->getSchedClass(SU);
for (TargetSchedModel::ProcResIter PI = SchedModel->getWriteProcResBegin(SC),
PE = SchedModel->getWriteProcResEnd(SC);
PI != PE; ++PI) {
@@ -563,27 +559,24 @@ unsigned CandidateHeuristics::getHWUICyclesForSU(SUnit *SU) {
return ReleaseAtCycle;
}
-void CandidateHeuristics::updateForScheduling(SUnit *SU) {
- HardwareUnitInfo *HWUI =
- getHWUIFromFlavor(classifyFlavor(*SU->getInstr(), *SII));
- assert(HWUI);
- HWUI->markScheduled(SU, getHWUICyclesForSU(SU));
+unsigned CandidateHeuristics::getHWUICyclesForSU(SUnit *SU) {
+ assert(SchedModel && SchedModel->hasInstrSchedModel());
+ MachineInstr *MI = SU->getInstr();
+ if (SII->isDS(*MI))
+ return SchedModel->computeInstrLatency(MI);
+ return getMaxBlockingCycles(DAG->getSchedClass(SU), MI);
}
unsigned CandidateHeuristics::getHWUICyclesForMI(MachineInstr *MI) {
assert(SchedModel && SchedModel->hasInstrSchedModel());
- // Loads and stores are not pipelined
- if (MI->mayLoadOrStore())
- return SchedModel->computeInstrLatency(MI, false);
+ return getMaxBlockingCycles(SchedModel->resolveSchedClass(MI), MI);
+}
- unsigned ReleaseAtCycle = 0;
- const MCSchedClassDesc *SC = SchedModel->resolveSchedClass(MI);
- for (TargetSchedModel::ProcResIter PI = SchedModel->getWriteProcResBegin(SC),
- PE = SchedModel->getWriteProcResEnd(SC);
- PI != PE; ++PI) {
- ReleaseAtCycle = std::max(ReleaseAtCycle, (unsigned)PI->ReleaseAtCycle);
- }
- return ReleaseAtCycle;
+void CandidateHeuristics::updateForScheduling(SUnit *SU) {
+ HardwareUnitInfo *HWUI =
+ getHWUIFromFlavor(classifyFlavor(*SU->getInstr(), *SII));
+ assert(HWUI);
+ HWUI->markScheduled(SU, getHWUICyclesForSU(SU));
}
void CandidateHeuristics::initialize(ScheduleDAGMI *SchedDAG,
@@ -800,15 +793,12 @@ bool CandidateHeuristics::tryEffectiveStall(
Costs.Ready = ReadyCycle > CurrCycle ? ReadyCycle - CurrCycle : 0;
Costs.Structural = getStructuralStallCycles(Zone, SU);
Costs.Latency = Zone.getLatencyStallCycles(SU);
+ unsigned CarriedLatency = CarriedLatencies.lookup_or(SU->getInstr(), 0);
+ Costs.Carried = CarriedLatency > CurrCycle ? CarriedLatency - CurrCycle : 0;
Costs.Buffer = getBufferFullStalls(SU);
- unsigned TryCarriedLatency =
- CarriedLatencies.contains(TryCand.SU->getInstr())
- ? CarriedLatencies[TryCand.SU->getInstr()]
- : 0;
- Costs.Carried =
- TryCarriedLatency > CurrCycle ? TryCarriedLatency - CurrCycle : 0;
+
Costs.Effective = std::max({Costs.Ready, Costs.Structural, Costs.Latency,
- Costs.Buffer, Costs.Carried});
+ Costs.Carried, Costs.Buffer});
return Costs;
};
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.h b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.h
index 5685091fc09c9..6e372f52d2069 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.h
@@ -209,6 +209,11 @@ class CandidateHeuristics {
/// heuristics.
void collectRegionSummary();
+ /// \returns the maximum blocking cycles according to the SchedModel for a
+ /// given MCSchedClassDesc \p SC
+ unsigned getMaxBlockingCycles(const MCSchedClassDesc *SC,
+ const MachineInstr *MI);
+
/// Compute the blocking cycles for the appropriate HardwareUnit given an \p
/// SU
unsigned getHWUICyclesForSU(SUnit *SU);
diff --git a/llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll b/llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll
index 0dc660143f2a8..096ec69c00029 100644
--- a/llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll
+++ b/llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll
@@ -286,73 +286,73 @@ define amdgpu_kernel void @ds_wmma_permute(ptr addrspace(3) %base, ptr addrspace
; COEXEC-NEXT: .LBB1_1: ; %loop
; COEXEC-NEXT: ; =>This Inner Loop Header: Depth=1
; COEXEC-NEXT: s_add_co_i32 s7, s0, s6
-; COEXEC-NEXT: s_add_co_i32 s8, s1, s6
-; COEXEC-NEXT: s_add_co_i32 s6, s6, s3
-; COEXEC-NEXT: s_and_b32 s9, s2, exec_lo
-; COEXEC-NEXT: s_cselect_b32 s9, 1, 0
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
-; COEXEC-NEXT: v_mov_b32_e32 v124, s7
-; COEXEC-NEXT: s_cmp_lg_u32 s9, 1
-; COEXEC-NEXT: v_mov_b32_e32 v156, s8
-; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v124
-; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v124 offset:64
-; COEXEC-NEXT: ds_load_tr16_b128 v[40:43], v156
-; COEXEC-NEXT: ds_load_tr16_b128 v[44:47], v156 offset:64
-; COEXEC-NEXT: ds_load_tr16_b128 v[52:55], v124 offset:320
-; COEXEC-NEXT: ds_load_tr16_b128 v[48:51], v124 offset:256
-; COEXEC-NEXT: ds_load_tr16_b128 v[56:59], v156 offset:256
-; COEXEC-NEXT: ds_load_tr16_b128 v[60:63], v156 offset:320
-; COEXEC-NEXT: ds_load_tr16_b128 v[68:71], v124 offset:576
-; COEXEC-NEXT: ds_load_tr16_b128 v[64:67], v124 offset:512
-; COEXEC-NEXT: ds_load_tr16_b128 v[76:79], v156 offset:576
-; COEXEC-NEXT: ds_load_tr16_b128 v[72:75], v156 offset:512
-; COEXEC-NEXT: ds_load_tr16_b128 v[84:87], v124 offset:832
-; COEXEC-NEXT: ds_load_tr16_b128 v[80:83], v124 offset:768
-; COEXEC-NEXT: ds_load_tr16_b128 v[92:95], v156 offset:832
-; COEXEC-NEXT: ds_load_tr16_b128 v[88:91], v156 offset:768
-; COEXEC-NEXT: ds_load_tr16_b128 v[96:99], v124 offset:128
-; COEXEC-NEXT: ds_load_tr16_b128 v[100:103], v124 offset:192
-; COEXEC-NEXT: ds_load_tr16_b128 v[104:107], v124 offset:384
-; COEXEC-NEXT: s_wait_dscnt 0xf
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[40:47], v[24:31]
-; COEXEC-NEXT: ds_load_tr16_b128 v[108:111], v124 offset:448
-; COEXEC-NEXT: ds_load_tr16_b128 v[112:115], v124 offset:640
-; COEXEC-NEXT: ds_load_tr16_b128 v[116:119], v124 offset:704
-; COEXEC-NEXT: ds_load_tr16_b128 v[120:123], v124 offset:896
-; COEXEC-NEXT: ds_load_tr16_b128 v[124:127], v124 offset:960
+; COEXEC-NEXT: v_mov_b32_e32 v92, s7
+; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v92 offset:64
+; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v92
+; COEXEC-NEXT: ds_load_tr16_b128 v[40:43], v92 offset:128
+; COEXEC-NEXT: ds_load_tr16_b128 v[44:47], v92 offset:192
+; COEXEC-NEXT: ds_load_tr16_b128 v[48:51], v92 offset:256
+; COEXEC-NEXT: ds_load_tr16_b128 v[52:55], v92 offset:320
+; COEXEC-NEXT: ds_load_tr16_b128 v[56:59], v92 offset:384
+; COEXEC-NEXT: ds_load_tr16_b128 v[60:63], v92 offset:448
+; COEXEC-NEXT: ds_load_tr16_b128 v[64:67], v92 offset:512
+; COEXEC-NEXT: ds_load_tr16_b128 v[68:71], v92 offset:576
+; COEXEC-NEXT: ds_load_tr16_b128 v[72:75], v92 offset:640
+; COEXEC-NEXT: ds_load_tr16_b128 v[76:79], v92 offset:704
+; COEXEC-NEXT: ds_load_tr16_b128 v[80:83], v92 offset:768
+; COEXEC-NEXT: ds_load_tr16_b128 v[84:87], v92 offset:832
+; COEXEC-NEXT: ds_load_tr16_b128 v[88:91], v92 offset:896
+; COEXEC-NEXT: ds_load_tr16_b128 v[92:95], v92 offset:960
+; COEXEC-NEXT: s_add_co_i32 s7, s1, s6
+; COEXEC-NEXT: s_add_co_i32 s6, s6, s3
+; COEXEC-NEXT: s_and_b32 s8, s2, exec_lo
+; COEXEC-NEXT: s_cselect_b32 s8, 1, 0
+; COEXEC-NEXT: v_mov_b32_e32 v156, s7
+; COEXEC-NEXT: s_cmp_lg_u32 s8, 1
+; COEXEC-NEXT: ds_load_tr16_b128 v[96:99], v156
+; COEXEC-NEXT: ds_load_tr16_b128 v[100:103], v156 offset:64
+; COEXEC-NEXT: ds_load_tr16_b128 v[108:111], v156 offset:320
+; COEXEC-NEXT: ds_load_tr16_b128 v[104:107], v156 offset:256
+; COEXEC-NEXT: ds_load_tr16_b128 v[116:119], v156 offset:576
+; COEXEC-NEXT: ds_load_tr16_b128 v[112:115], v156 offset:512
+; COEXEC-NEXT: ds_load_tr16_b128 v[124:127], v156 offset:832
+; COEXEC-NEXT: ds_load_tr16_b128 v[120:123], v156 offset:768
; COEXEC-NEXT: ds_load_tr16_b128 v[128:131], v156 offset:128
; COEXEC-NEXT: ds_load_tr16_b128 v[132:135], v156 offset:192
-; COEXEC-NEXT: s_wait_dscnt 0x12
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[56:63], v[16:23]
; COEXEC-NEXT: ds_load_tr16_b128 v[136:139], v156 offset:384
; COEXEC-NEXT: ds_load_tr16_b128 v[140:143], v156 offset:448
; COEXEC-NEXT: ds_load_tr16_b128 v[144:147], v156 offset:640
; COEXEC-NEXT: ds_load_tr16_b128 v[148:151], v156 offset:704
; COEXEC-NEXT: ds_load_tr16_b128 v[152:155], v156 offset:896
; COEXEC-NEXT: ds_load_tr16_b128 v[156:159], v156 offset:960
-; COEXEC-NEXT: s_wait_dscnt 0x14
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[72:79], v[8:15]
-; COEXEC-NEXT: s_wait_dscnt 0x10
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[80:87], v[88:95], v[0:7]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[40:47], v[24:31]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[56:63], v[16:23]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[72:79], v[8:15]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[80:87], v[88:95], v[0:7]
+; COEXEC-NEXT: s_wait_dscnt 0xe
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[96:103], v[24:31]
+; COEXEC-NEXT: s_wait_dscnt 0xc
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[104:111], v[16:23]
+; COEXEC-NEXT: s_wait_dscnt 0xa
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[112:119], v[8:15]
+; COEXEC-NEXT: s_wait_dscnt 0x8
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[80:87], v[120:127], v[0:7]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[96:103], v[24:31]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[104:111], v[16:23]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[112:119], v[8:15]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[80:87], v[120:127], v[0:7]
; COEXEC-NEXT: s_wait_dscnt 0x6
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[96:103], v[128:135], v[24:31]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[128:135], v[24:31]
; COEXEC-NEXT: s_wait_dscnt 0x4
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[104:111], v[136:143], v[16:23]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[56:63], v[136:143], v[16:23]
; COEXEC-NEXT: s_wait_dscnt 0x2
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[112:119], v[144:151], v[8:15]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[72:79], v[144:151], v[8:15]
; COEXEC-NEXT: s_wait_dscnt 0x0
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[120:127], v[152:159], v[0:7]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[96:103], v[128:135], v[24:31]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[104:111], v[136:143], v[16:23]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[112:119], v[144:151], v[8:15]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[120:127], v[152:159], v[0:7]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[88:95], v[152:159], v[0:7]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[128:135], v[24:31]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[56:63], v[136:143], v[16:23]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[72:79], v[144:151], v[8:15]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[88:95], v[152:159], v[0:7]
; COEXEC-NEXT: s_cbranch_scc1 .LBB1_1
; COEXEC-NEXT: ; %bb.2: ; %end
; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
@@ -645,23 +645,21 @@ define amdgpu_kernel void @ds_wmma_block_carried(ptr addrspace(3) %base, ptr add
; COEXEC-NEXT: s_add_co_i32 s0, s3, s1
; COEXEC-NEXT: v_mov_b32_e32 v31, v0
; COEXEC-NEXT: s_xor_b32 s2, s2, -1
-; COEXEC-NEXT: v_mov_b32_e32 v52, s3
-; COEXEC-NEXT: ds_load_tr16_b128 v[40:43], v52
-; COEXEC-NEXT: ds_load_tr16_b128 v[44:47], v52 offset:64
-; COEXEC-NEXT: ds_load_tr16_b128 v[56:59], v52 offset:128
-; COEXEC-NEXT: ds_load_tr16_b128 v[60:63], v52 offset:192
-; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v52 offset:256
-; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v52 offset:320
-; COEXEC-NEXT: ds_load_tr16_b128 v[48:51], v52 offset:384
-; COEXEC-NEXT: ds_load_tr16_b128 v[52:55], v52 offset:448
+; COEXEC-NEXT: v_mov_b32_e32 v60, s3
+; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v60
+; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v60 offset:64
+; COEXEC-NEXT: ds_load_tr16_b128 v[40:43], v60 offset:128
+; COEXEC-NEXT: ds_load_tr16_b128 v[44:47], v60 offset:192
+; COEXEC-NEXT: ds_load_tr16_b128 v[48:51], v60 offset:256
+; COEXEC-NEXT: ds_load_tr16_b128 v[52:55], v60 offset:320
+; COEXEC-NEXT: ds_load_tr16_b128 v[56:59], v60 offset:384
+; COEXEC-NEXT: ds_load_tr16_b128 v[60:63], v60 offset:448
; COEXEC-NEXT: s_wait_dscnt 0x0
; COEXEC-NEXT: .LBB2_1: ; %loop
; COEXEC-NEXT: ; =>This Inner Loop Header: Depth=1
-; COEXEC-NEXT: s_wait_dscnt 0x4
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[56:63], v[24:31]
-; COEXEC-NEXT: s_wait_dscnt 0x0
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[32:39], v[48:55], v[16:23]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[56:63], v[24:31]
+; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_mov_b64_e32 v[70:71], s[14:15]
; COEXEC-NEXT: v_mov_b64_e32 v[68:69], s[12:13]
@@ -671,24 +669,29 @@ define amdgpu_kernel void @ds_wmma_block_carried(ptr addrspace(3) %base, ptr add
; COEXEC-NEXT: s_add_co_i32 s0, s0, s1
; COEXEC-NEXT: s_and_b32 s3, s2, exec_lo
; COEXEC-NEXT: s_cselect_b32 s3, 1, 0
-; COEXEC-NEXT: ds_load_tr16_b128 v[40:43], v72
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[64:71], v[8:15]
-; COEXEC-NEXT: ds_load_tr16_b128 v[44:47], v72 offset:64
-; COEXEC-NEXT: ds_load_tr16_b128 v[56:59], v72 offset:128
-; COEXEC-NEXT: ds_load_tr16_b128 v[60:63], v72 offset:192
; COEXEC-NEXT: s_cmp_lg_u32 s3, 1
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[64:71], v[64:71], v[0:7]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[32:39], v[48:55], v[16:23]
-; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v72 offset:256
-; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v72 offset:320
-; COEXEC-NEXT: ds_load_tr16_b128 v[48:51], v72 offset:384
-; COEXEC-NEXT: ds_load_tr16_b128 v[52:55], v72 offset:448
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[64:71], v[8:15]
+; COEXEC-NEXT: s_wait_dscnt 0x4
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[40:47], v[24:31]
+; COEXEC-NEXT: s_wait_dscnt 0x0
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[56:63], v[16:23]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[64:71], v[64:71], v[0:7]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[40:47], v[24:31]
+; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v72
+; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v72 offset:64
+; COEXEC-NEXT: ds_load_tr16_b128 v[40:43], v72 offset:128
+; COEXEC-NEXT: ds_load_tr16_b128 v[44:47], v72 offset:192
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[56:63], v[16:23]
+; COEXEC-NEXT: ds_load_tr16_b128 v[48:51], v72 offset:256
+; COEXEC-NEXT: ds_load_tr16_b128 v[52:55], v72 offset:320
+; COEXEC-NEXT: ds_load_tr16_b128 v[56:59], v72 offset:384
+; COEXEC-NEXT: ds_load_tr16_b128 v[60:63], v72 offset:448
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[64:71], v[64:71], v[0:7]
; COEXEC-NEXT: s_cbranch_scc1 .LBB2_1
; COEXEC-NEXT: ; %bb.2: ; %end
; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
-; COEXEC-NEXT: s_wait_dscnt 0x3
+; COEXEC-NEXT: s_wait_dscnt 0x7
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_mov_b32_e32 v32, 0
@@ -932,46 +935,44 @@ define amdgpu_kernel void @ds_wmma_loop_carried(ptr addrspace(3) %base, ptr addr
; COEXEC-NEXT: s_bitcmp1_b32 s0, 0
; COEXEC-NEXT: s_cselect_b32 s3, -1, 0
; COEXEC-NEXT: s_add_co_i32 s0, s2, s1
-; COEXEC-NEXT: ds_load_tr16_b128 v[8:11], v40
-; COEXEC-NEXT: ds_load_tr16_b128 v[12:15], v40 offset:64
-; COEXEC-NEXT: ds_load_tr16_b128 v[4:7], v40 offset:320
-; COEXEC-NEXT: ds_load_tr16_b128 v[0:3], v40 offset:256
+; COEXEC-NEXT: ds_load_tr16_b128 v[0:3], v40
+; COEXEC-NEXT: ds_load_tr16_b128 v[4:7], v40 offset:64
+; COEXEC-NEXT: ds_load_tr16_b128 v[12:15], v40 offset:320
+; COEXEC-NEXT: ds_load_tr16_b128 v[8:11], v40 offset:256
; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v40 offset:576
; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v40 offset:512
; COEXEC-NEXT: ds_load_tr16_b128 v[76:79], v40 offset:832
; COEXEC-NEXT: ds_load_tr16_b128 v[72:75], v40 offset:768
-; COEXEC-NEXT: ds_load_tr16_b128 v[24:27], v40 offset:128
-; COEXEC-NEXT: ds_load_tr16_b128 v[28:31], v40 offset:192
-; COEXEC-NEXT: ds_load_tr16_b128 v[16:19], v40 offset:384
-; COEXEC-NEXT: ds_load_tr16_b128 v[20:23], v40 offset:448
+; COEXEC-NEXT: ds_load_tr16_b128 v[16:19], v40 offset:128
+; COEXEC-NEXT: ds_load_tr16_b128 v[20:23], v40 offset:192
+; COEXEC-NEXT: ds_load_tr16_b128 v[24:27], v40 offset:384
+; COEXEC-NEXT: ds_load_tr16_b128 v[28:31], v40 offset:448
; COEXEC-NEXT: ds_load_tr16_b128 v[80:83], v40 offset:640
; COEXEC-NEXT: ds_load_tr16_b128 v[84:87], v40 offset:704
; COEXEC-NEXT: ds_load_tr16_b128 v[88:91], v40 offset:896
; COEXEC-NEXT: ds_load_tr16_b128 v[92:95], v40 offset:960
; COEXEC-NEXT: s_xor_b32 s2, s3, -1
; COEXEC-NEXT: s_wait_dscnt 0xe
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[8:15], v[64:71], 0
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[0:7], v[64:71], 0
; COEXEC-NEXT: s_wait_dscnt 0xc
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[48:55], v[0:7], v[64:71], 0
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[48:55], v[8:15], v[64:71], 0
; COEXEC-NEXT: s_wait_dscnt 0xa
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[40:47], v[32:39], v[64:71], 0
; COEXEC-NEXT: s_wait_dscnt 0x8
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[72:79], v[64:71], 0
; COEXEC-NEXT: s_wait_dscnt 0x6
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[24:31], v[64:71], v[56:63]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[16:23], v[64:71], v[56:63]
; COEXEC-NEXT: s_wait_dscnt 0x4
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[48:55], v[16:23], v[64:71], v[48:55]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[48:55], v[24:31], v[64:71], v[48:55]
; COEXEC-NEXT: s_wait_dscnt 0x2
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[40:47], v[80:87], v[64:71], v[40:47]
; COEXEC-NEXT: s_wait_dscnt 0x0
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[88:95], v[64:71], v[32:39]
; COEXEC-NEXT: .LBB3_1: ; %loop
; COEXEC-NEXT: ; =>This Inner Loop Header: Depth=1
-; COEXEC-NEXT: s_wait_dscnt 0x4
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[8:15], v[24:31], v[56:63]
-; COEXEC-NEXT: s_wait_dscnt 0x0
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[48:55], v[0:7], v[16:23], v[48:55]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[8:15], v[24:31], v[56:63]
+; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_mov_b64_e32 v[70:71], s[14:15]
; COEXEC-NEXT: v_mov_b64_e32 v[68:69], s[12:13]
@@ -981,24 +982,29 @@ define amdgpu_kernel void @ds_wmma_loop_carried(ptr addrspace(3) %base, ptr addr
; COEXEC-NEXT: s_add_co_i32 s0, s0, s1
; COEXEC-NEXT: s_and_b32 s3, s2, exec_lo
; COEXEC-NEXT: s_cselect_b32 s3, 1, 0
-; COEXEC-NEXT: ds_load_tr16_b128 v[8:11], v72
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[40:47], v[64:71], v[64:71], v[40:47]
-; COEXEC-NEXT: ds_load_tr16_b128 v[12:15], v72 offset:64
-; COEXEC-NEXT: ds_load_tr16_b128 v[24:27], v72 offset:128
-; COEXEC-NEXT: ds_load_tr16_b128 v[28:31], v72 offset:192
; COEXEC-NEXT: s_cmp_lg_u32 s3, 1
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[64:71], v[64:71], v[32:39]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[48:55], v[0:7], v[16:23], v[48:55]
-; COEXEC-NEXT: ds_load_tr16_b128 v[0:3], v72 offset:256
-; COEXEC-NEXT: ds_load_tr16_b128 v[4:7], v72 offset:320
-; COEXEC-NEXT: ds_load_tr16_b128 v[16:19], v72 offset:384
-; COEXEC-NEXT: ds_load_tr16_b128 v[20:23], v72 offset:448
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[40:47], v[64:71], v[64:71], v[40:47]
+; COEXEC-NEXT: s_wait_dscnt 0x4
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[0:7], v[16:23], v[56:63]
+; COEXEC-NEXT: s_wait_dscnt 0x0
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[48:55], v[8:15], v[24:31], v[48:55]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[64:71], v[64:71], v[32:39]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[0:7], v[16:23], v[56:63]
+; COEXEC-NEXT: ds_load_tr16_b128 v[0:3], v72
+; COEXEC-NEXT: ds_load_tr16_b128 v[4:7], v72 offset:64
+; COEXEC-NEXT: ds_load_tr16_b128 v[16:19], v72 offset:128
+; COEXEC-NEXT: ds_load_tr16_b128 v[20:23], v72 offset:192
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[48:55], v[8:15], v[24:31], v[48:55]
+; COEXEC-NEXT: ds_load_tr16_b128 v[8:11], v72 offset:256
+; COEXEC-NEXT: ds_load_tr16_b128 v[12:15], v72 offset:320
+; COEXEC-NEXT: ds_load_tr16_b128 v[24:27], v72 offset:384
+; COEXEC-NEXT: ds_load_tr16_b128 v[28:31], v72 offset:448
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[64:71], v[64:71], v[32:39]
; COEXEC-NEXT: s_cbranch_scc1 .LBB3_1
; COEXEC-NEXT: ; %bb.2: ; %end
; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
-; COEXEC-NEXT: s_wait_dscnt 0x3
+; COEXEC-NEXT: s_wait_dscnt 0x7
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_mov_b32_e32 v0, 0
>From 81337caf499c2b5fb8dfd07fb20896aa758ff7ec Mon Sep 17 00:00:00 2001
From: Jeffrey Byrnes <Jeffrey.Byrnes at amd.com>
Date: Thu, 28 May 2026 11:46:34 -0700
Subject: [PATCH 03/10] Merge conflict
Change-Id: I24f471688f9d0604b45e95a4fa4da85fb0d9ed76
---
.../AMDGPU/AMDGPUCoExecSchedStrategy.cpp | 35 +++++++++----------
1 file changed, 17 insertions(+), 18 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
index 3311a4432a102..320aace393ebb 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
@@ -756,13 +756,23 @@ bool CandidateHeuristics::tryEffectiveStall(
GenericSchedulerBase::SchedCandidate &Cand,
GenericSchedulerBase::SchedCandidate &TryCand, SchedBoundary &Zone) {
+ // Treat structural and latency stalls as a single scheduling cost for the
+ // current cycle.
+ struct StallCosts {
+ unsigned Ready = 0;
+ unsigned Structural = 0;
+ unsigned Latency = 0;
+ unsigned Effective = 0;
+ unsigned Carried = 0;
+ unsigned Buffer = 0;
+ };
+
auto getBufferFullStalls = [this, &Zone](SUnit *SU) -> unsigned {
InstructionFlavor Flavor = classifyFlavor(
*SU->getInstr(), *static_cast<const SIInstrInfo *>(DAG->TII));
HardwareUnitInfo *HWUI = getHWUIFromFlavor(Flavor);
- // A BufferSize of 0 means "unlimited" buffer, thus we will never fill it.
- if (HWUI->getBufferSize() == 0)
+ if (HWUI->getBufferSize() <= 1)
return 0;
// getBufferAvailableCycle assumes top-down scheduling.
@@ -775,17 +785,6 @@ bool CandidateHeuristics::tryEffectiveStall(
return BufferReadyCycle - CurrCycle;
};
- // Treat structural and latency stalls as a single scheduling cost for the
- // current cycle.
- struct StallCosts {
- unsigned Ready = 0;
- unsigned Structural = 0;
- unsigned Latency = 0;
- unsigned Effective = 0;
- unsigned Buffer = 0;
- unsigned Carried = 0;
- };
-
unsigned CurrCycle = Zone.getCurrCycle();
auto GetStallCosts = [&](SUnit *SU) {
unsigned ReadyCycle = Zone.isTop() ? SU->TopReadyCycle : SU->BotReadyCycle;
@@ -808,12 +807,12 @@ bool CandidateHeuristics::tryEffectiveStall(
LLVM_DEBUG(if (TryCosts.Effective || CandCosts.Effective) {
dbgs() << "Effective stalls: try=" << TryCosts.Effective
<< " (ready=" << TryCosts.Ready << ", struct=" << TryCosts.Structural
- << ", lat=" << TryCosts.Latency << ", buffer=" << TryCosts.Buffer
- << ", carried=" << TryCosts.Carried
- << ") cand=" << CandCosts.Effective << " (ready=" << CandCosts.Ready
+ << ", lat=" << TryCosts.Latency << ", carried=" << TryCosts.Carried
+ << ", buffer=" << TryCosts.Buffer << ") cand=" << CandCosts.Effective
+ << " (ready=" << CandCosts.Ready
<< ", struct=" << CandCosts.Structural
- << ", lat=" << CandCosts.Latency << ", buffer=" << CandCosts.Buffer
- << ", carried=" << CandCosts.Carried << ")\n";
+ << ", lat=" << CandCosts.Latency << ", carried=" << CandCosts.Carried
+ << ", buffer=" << CandCosts.Buffer << ")\n";
});
return tryLess(TryCosts.Effective, CandCosts.Effective, TryCand, Cand,
>From a8512fff08e0a34d02695e07dcc4e1b99e32845b Mon Sep 17 00:00:00 2001
From: Jeffrey Byrnes <Jeffrey.Byrnes at amd.com>
Date: Wed, 15 Apr 2026 14:23:08 -0700
Subject: [PATCH 04/10] Address comments from
https://github.com/llvm/llvm-project/pull/187413
Change-Id: I0fad2fc504d42b72a664f78823e63c63ba2e4045
---
llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp | 10 +++-------
1 file changed, 3 insertions(+), 7 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
index 320aace393ebb..a0b46925dd840 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
@@ -607,16 +607,12 @@ void CandidateHeuristics::initialize(ScheduleDAGMI *SchedDAG,
unsigned CandidateHeuristics::getCarriedLatency(SUnit *SU) {
MachineInstr *MI = SU->getInstr();
unsigned CarriedLatency = 0;
- for (auto &Op : MI->operands()) {
- if (!Op.isReg())
- continue;
- if (!Op.isUse())
- continue;
+ for (MachineOperand &Op : MI->all_uses()) {
auto Reg = Op.getReg();
if (!Reg.isVirtual())
continue;
- for (auto &Def : DAG->MRI.def_instructions(Reg)) {
+ for (MachineInstr &Def : DAG->MRI.def_instructions(Reg)) {
// We don't have the proper modelling to accurately measure all carried
// latency. Just try to measure carried latency for long latency loads to
// avoid long stalls.
@@ -628,7 +624,7 @@ unsigned CandidateHeuristics::getCarriedLatency(SUnit *SU) {
// Load is carried across block
if (Def.getParent() != MI->getParent()) {
bool FoundUseInDefBlock = false;
- for (auto &Use : DAG->MRI.use_nodbg_instructions(Reg)) {
+ for (MachineInstr &Use : DAG->MRI.use_nodbg_instructions(Reg)) {
if (Use.getParent() != Def.getParent())
continue;
>From f6444a6eaa1d596ddc4fdb7f972b19e9b48c3b4e Mon Sep 17 00:00:00 2001
From: Jeffrey Byrnes <Jeffrey.Byrnes at amd.com>
Date: Fri, 24 Apr 2026 09:54:15 -0700
Subject: [PATCH 05/10] Remove unused function
Change-Id: I9f2de1497f793d2848dedaf645e21e07a4ba82d6
---
llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
index a0b46925dd840..8aa16cbacfdef 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
@@ -768,7 +768,8 @@ bool CandidateHeuristics::tryEffectiveStall(
*SU->getInstr(), *static_cast<const SIInstrInfo *>(DAG->TII));
HardwareUnitInfo *HWUI = getHWUIFromFlavor(Flavor);
- if (HWUI->getBufferSize() <= 1)
+ // A BufferSize of 0 means "unlimited" buffer, thus we will never fill it.
+ if (HWUI->getBufferSize() == 0)
return 0;
// getBufferAvailableCycle assumes top-down scheduling.
>From 399aa9150669238c30754f022df08badf2b42b05 Mon Sep 17 00:00:00 2001
From: Jeffrey Byrnes <Jeffrey.Byrnes at amd.com>
Date: Fri, 24 Apr 2026 10:21:50 -0700
Subject: [PATCH 06/10] Use static_cast
Change-Id: Ibec2cf245d5ac213ef0cc4292ba80cb983a58692
---
.../AMDGPU/AMDGPUCoExecSchedStrategy.cpp | 19 ++++++++++++-------
1 file changed, 12 insertions(+), 7 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
index 8aa16cbacfdef..fc85df03a81f8 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
@@ -554,7 +554,8 @@ unsigned CandidateHeuristics::getMaxBlockingCycles(const MCSchedClassDesc *SC,
for (TargetSchedModel::ProcResIter PI = SchedModel->getWriteProcResBegin(SC),
PE = SchedModel->getWriteProcResEnd(SC);
PI != PE; ++PI) {
- ReleaseAtCycle = std::max(ReleaseAtCycle, (unsigned)PI->ReleaseAtCycle);
+ ReleaseAtCycle =
+ std::max(ReleaseAtCycle, static_cast<unsigned>(PI->ReleaseAtCycle));
}
return ReleaseAtCycle;
}
@@ -589,17 +590,21 @@ void CandidateHeuristics::initialize(ScheduleDAGMI *SchedDAG,
SRI = static_cast<const SIRegisterInfo *>(TRI);
SII = static_cast<const SIInstrInfo *>(DAG->TII);
- HWUInfo.resize((int)InstructionFlavor::NUM_FLAVORS);
+ HWUInfo.resize(static_cast<int>(InstructionFlavor::NUM_FLAVORS));
for (unsigned I = 0; I < HWUInfo.size(); I++) {
HWUInfo[I].reset();
HWUInfo[I].setType(I);
}
- HWUInfo[(int)InstructionFlavor::WMMA].setProducesCoexecWindow(true);
- HWUInfo[(int)InstructionFlavor::MultiCycleVALU].setProducesCoexecWindow(true);
- HWUInfo[(int)InstructionFlavor::TRANS].setProducesCoexecWindow(true);
- HWUInfo[(int)InstructionFlavor::DS].setBufferSize(DefaultBufferSizes::DS);
+ HWUInfo[static_cast<int>(InstructionFlavor::WMMA)].setProducesCoexecWindow(
+ true);
+ HWUInfo[static_cast<int>(InstructionFlavor::MultiCycleVALU)]
+ .setProducesCoexecWindow(true);
+ HWUInfo[static_cast<int>(InstructionFlavor::TRANS)].setProducesCoexecWindow(
+ true);
+ HWUInfo[static_cast<int>(InstructionFlavor::DS)].setBufferSize(
+ DefaultBufferSizes::DS);
collectRegionSummary();
}
@@ -664,7 +669,7 @@ void CandidateHeuristics::collectRegionSummary() {
for (auto &SU : DAG->SUnits) {
MachineInstr *MI = SU.getInstr();
const InstructionFlavor Flavor = classifyFlavor(*MI, *SII);
- HWUInfo[(int)(Flavor)].insert(&SU, getHWUICyclesForSU(&SU));
+ HWUInfo[static_cast<int>(Flavor)].insert(&SU, getHWUICyclesForSU(&SU));
unsigned CarriedLatency = getCarriedLatency(&SU);
if (CarriedLatency)
CarriedLatencies[MI] = CarriedLatency;
>From f88e08dc709eae4974ea8394348f16080dc8fecd Mon Sep 17 00:00:00 2001
From: Jeffrey Byrnes <Jeffrey.Byrnes at amd.com>
Date: Thu, 28 May 2026 13:51:18 -0700
Subject: [PATCH 07/10] Add support for fence in carried latency
Change-Id: I340660e56e74b559594021e7f0f6004bbfc1c3bf
---
.../AMDGPU/AMDGPUCoExecSchedStrategy.cpp | 42 +++++++++++++++++++
1 file changed, 42 insertions(+)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
index fc85df03a81f8..994646d2254c6 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
@@ -20,6 +20,22 @@ using namespace llvm;
using namespace llvm::AMDGPU;
#define DEBUG_TYPE "machine-scheduler"
+namespace {
+enum class CarriedLatency { Off, Fence, All };
+} // namespace
+
+static cl::opt<CarriedLatency> BlockCarriedLatency(
+ "amdgpu-block-carried-latency", cl::Hidden,
+ cl::init(CarriedLatency::Off),
+ cl::desc("Prioritize HardwareUnits with non-zero exposed cycles in the "
+ "coexec scheduler's critical-resource sort."),
+ cl::values(
+ clEnumValN(CarriedLatency::Off, "off",
+ "Disabled- do not pad latency."),
+ clEnumValN(CarriedLatency::Fence, "fence",
+ "Only pad latency for memory fence (e.g. those surrounding barrier_signal/wait)."),
+ clEnumValN(CarriedLatency::All, "all",
+ "Pad latency for any SU with an incoming ds_load dependency.")));
namespace {
@@ -610,7 +626,33 @@ void CandidateHeuristics::initialize(ScheduleDAGMI *SchedDAG,
}
unsigned CandidateHeuristics::getCarriedLatency(SUnit *SU) {
+ if (BlockCarriedLatency == CarriedLatency::Off)
+ return 0;
+
MachineInstr *MI = SU->getInstr();
+ const InstructionFlavor Flavor = classifyFlavor(*MI, *SII);
+ if (Flavor == InstructionFlavor::Fence) {
+ MachineBasicBlock *MBB = MI->getParent();
+ // Check if we have DS instruciton after a fence in any of the predecessor
+ // blocks, if so, this fence instruction has carried latency.
+ for (auto PredMBB : MBB->predecessors()) {
+ auto I = PredMBB->rbegin();
+ auto E = PredMBB->rend();
+ for (; I != E; I++) {
+ const InstructionFlavor ItFlavor = classifyFlavor(*I, *SII);
+ if (ItFlavor == InstructionFlavor::Fence)
+ break;
+
+ // Found carried latency.
+ if (ItFlavor == InstructionFlavor::DS)
+ return getHWUICyclesForMI(&*I);
+ }
+ }
+ }
+
+ if (BlockCarriedLatency == CarriedLatency::Fence)
+ return 0;
+
unsigned CarriedLatency = 0;
for (MachineOperand &Op : MI->all_uses()) {
auto Reg = Op.getReg();
>From 3a57d53c5ed84ecb95f34207ebf7955064d5936b Mon Sep 17 00:00:00 2001
From: Jeffrey Byrnes <Jeffrey.Byrnes at amd.com>
Date: Thu, 28 May 2026 14:26:03 -0700
Subject: [PATCH 08/10] Add test for block-carried modes
Change-Id: I4c521c05ec401581726ac040daadab6aceb8bd20
---
.../AMDGPU/AMDGPUCoExecSchedStrategy.cpp | 14 +-
.../AMDGPU/coexec-block-carried-latency.mir | 201 ++++++++++++++++++
2 files changed, 208 insertions(+), 7 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/coexec-block-carried-latency.mir
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
index 994646d2254c6..2e0a85ccb5fd5 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
@@ -25,17 +25,17 @@ enum class CarriedLatency { Off, Fence, All };
} // namespace
static cl::opt<CarriedLatency> BlockCarriedLatency(
- "amdgpu-block-carried-latency", cl::Hidden,
- cl::init(CarriedLatency::Off),
+ "amdgpu-block-carried-latency", cl::Hidden, cl::init(CarriedLatency::Off),
cl::desc("Prioritize HardwareUnits with non-zero exposed cycles in the "
"coexec scheduler's critical-resource sort."),
cl::values(
- clEnumValN(CarriedLatency::Off, "off",
- "Disabled- do not pad latency."),
+ clEnumValN(CarriedLatency::Off, "off", "Disabled- do not pad latency."),
clEnumValN(CarriedLatency::Fence, "fence",
- "Only pad latency for memory fence (e.g. those surrounding barrier_signal/wait)."),
- clEnumValN(CarriedLatency::All, "all",
- "Pad latency for any SU with an incoming ds_load dependency.")));
+ "Only pad latency for memory fence (e.g. those surrounding "
+ "barrier_signal/wait)."),
+ clEnumValN(
+ CarriedLatency::All, "all",
+ "Pad latency for any SU with an incoming ds_load dependency.")));
namespace {
diff --git a/llvm/test/CodeGen/AMDGPU/coexec-block-carried-latency.mir b/llvm/test/CodeGen/AMDGPU/coexec-block-carried-latency.mir
new file mode 100644
index 0000000000000..289a59fd1f032
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/coexec-block-carried-latency.mir
@@ -0,0 +1,201 @@
+# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 6
+# RUN: llc -mtriple=amdgcn -mcpu=gfx1250 -amdgpu-sched-strategy=coexec --run-pass=machine-scheduler -verify-misched %s -o - | FileCheck -check-prefix=DEFAULT %s
+# RUN: llc -mtriple=amdgcn -mcpu=gfx1250 -amdgpu-sched-strategy=coexec -amdgpu-block-carried-latency=fence --run-pass=machine-scheduler -verify-misched %s -o - | FileCheck -check-prefix=FENCE %s
+# RUN: llc -mtriple=amdgcn -mcpu=gfx1250 -amdgpu-sched-strategy=coexec -amdgpu-block-carried-latency=all --run-pass=machine-scheduler -verify-misched %s -o - | FileCheck -check-prefix=ALL %s
+
+--- |
+ define void @test-block-carried-latency() #0 { ret void }
+
+ attributes #0 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="1,128" }
+...
+
+---
+name: test-block-carried-latency
+tracksRegLiveness: true
+body: |
+ ; DEFAULT-LABEL: name: test-block-carried-latency
+ ; DEFAULT: bb.0:
+ ; DEFAULT-NEXT: successors: %bb.1(0x80000000)
+ ; DEFAULT-NEXT: liveins: $vgpr0, $sgpr0, $sgpr1
+ ; DEFAULT-NEXT: {{ $}}
+ ; DEFAULT-NEXT: [[DS_READ_B32_gfx9_:%[0-9]+]]:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 0, 0, implicit $exec
+ ; DEFAULT-NEXT: [[DS_READ_B32_gfx9_1:%[0-9]+]]:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 4, 0, implicit $exec
+ ; DEFAULT-NEXT: [[DS_READ_B32_gfx9_2:%[0-9]+]]:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 8, 0, implicit $exec
+ ; DEFAULT-NEXT: [[DS_READ_B32_gfx9_3:%[0-9]+]]:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 12, 0, implicit $exec
+ ; DEFAULT-NEXT: [[DEF:%[0-9]+]]:vreg_512_align2 = IMPLICIT_DEF
+ ; DEFAULT-NEXT: [[DEF1:%[0-9]+]]:vreg_512_align2 = IMPLICIT_DEF
+ ; DEFAULT-NEXT: [[DEF2:%[0-9]+]]:vreg_256_align2 = IMPLICIT_DEF
+ ; DEFAULT-NEXT: [[DEF3:%[0-9]+]]:vreg_512_align2 = IMPLICIT_DEF
+ ; DEFAULT-NEXT: [[DEF4:%[0-9]+]]:vreg_512_align2 = IMPLICIT_DEF
+ ; DEFAULT-NEXT: [[DEF5:%[0-9]+]]:vreg_256_align2 = IMPLICIT_DEF
+ ; DEFAULT-NEXT: S_BRANCH %bb.1
+ ; DEFAULT-NEXT: {{ $}}
+ ; DEFAULT-NEXT: bb.1:
+ ; DEFAULT-NEXT: liveins: $vgpr0, $sgpr0, $sgpr1
+ ; DEFAULT-NEXT: {{ $}}
+ ; DEFAULT-NEXT: early-clobber %10:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF]], [[DEF1]], 0, [[DEF2]], [[DS_READ_B32_gfx9_]], [[DS_READ_B32_gfx9_1]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead [[S_ADD_I32_:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; DEFAULT-NEXT: dead [[S_ADD_I32_1:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; DEFAULT-NEXT: dead [[S_ADD_I32_2:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; DEFAULT-NEXT: dead [[S_ADD_I32_3:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead [[S_ADD_I32_4:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_1:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead [[S_ADD_I32_5:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_2:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead [[S_ADD_I32_6:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_3:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead [[S_ADD_I32_7:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_4:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead [[S_ADD_I32_8:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_5:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; DEFAULT-NEXT: early-clobber %11:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead [[S_ADD_I32_9:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; DEFAULT-NEXT: ATOMIC_FENCE 5, 2
+ ; DEFAULT-NEXT: dead early-clobber %28:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead early-clobber %29:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead early-clobber %30:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead early-clobber %31:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead early-clobber %32:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; DEFAULT-NEXT: S_ENDPGM 0, implicit %10, implicit %11
+ ;
+ ; FENCE-LABEL: name: test-block-carried-latency
+ ; FENCE: bb.0:
+ ; FENCE-NEXT: successors: %bb.1(0x80000000)
+ ; FENCE-NEXT: liveins: $vgpr0, $sgpr0, $sgpr1
+ ; FENCE-NEXT: {{ $}}
+ ; FENCE-NEXT: [[DS_READ_B32_gfx9_:%[0-9]+]]:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 0, 0, implicit $exec
+ ; FENCE-NEXT: [[DS_READ_B32_gfx9_1:%[0-9]+]]:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 4, 0, implicit $exec
+ ; FENCE-NEXT: [[DS_READ_B32_gfx9_2:%[0-9]+]]:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 8, 0, implicit $exec
+ ; FENCE-NEXT: [[DS_READ_B32_gfx9_3:%[0-9]+]]:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 12, 0, implicit $exec
+ ; FENCE-NEXT: [[DEF:%[0-9]+]]:vreg_512_align2 = IMPLICIT_DEF
+ ; FENCE-NEXT: [[DEF1:%[0-9]+]]:vreg_512_align2 = IMPLICIT_DEF
+ ; FENCE-NEXT: [[DEF2:%[0-9]+]]:vreg_256_align2 = IMPLICIT_DEF
+ ; FENCE-NEXT: [[DEF3:%[0-9]+]]:vreg_512_align2 = IMPLICIT_DEF
+ ; FENCE-NEXT: [[DEF4:%[0-9]+]]:vreg_512_align2 = IMPLICIT_DEF
+ ; FENCE-NEXT: [[DEF5:%[0-9]+]]:vreg_256_align2 = IMPLICIT_DEF
+ ; FENCE-NEXT: S_BRANCH %bb.1
+ ; FENCE-NEXT: {{ $}}
+ ; FENCE-NEXT: bb.1:
+ ; FENCE-NEXT: liveins: $vgpr0, $sgpr0, $sgpr1
+ ; FENCE-NEXT: {{ $}}
+ ; FENCE-NEXT: early-clobber %10:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF]], [[DEF1]], 0, [[DEF2]], [[DS_READ_B32_gfx9_]], [[DS_READ_B32_gfx9_1]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; FENCE-NEXT: dead [[S_ADD_I32_:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; FENCE-NEXT: dead [[S_ADD_I32_1:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; FENCE-NEXT: dead [[S_ADD_I32_2:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; FENCE-NEXT: dead [[S_ADD_I32_3:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; FENCE-NEXT: dead [[V_ADD_U32_e64_:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; FENCE-NEXT: dead [[S_ADD_I32_4:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; FENCE-NEXT: dead [[V_ADD_U32_e64_1:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; FENCE-NEXT: dead [[S_ADD_I32_5:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; FENCE-NEXT: dead [[V_ADD_U32_e64_2:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; FENCE-NEXT: dead [[S_ADD_I32_6:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; FENCE-NEXT: dead [[V_ADD_U32_e64_3:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; FENCE-NEXT: dead [[S_ADD_I32_7:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; FENCE-NEXT: dead [[V_ADD_U32_e64_4:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; FENCE-NEXT: dead [[S_ADD_I32_8:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; FENCE-NEXT: dead [[V_ADD_U32_e64_5:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; FENCE-NEXT: early-clobber %11:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; FENCE-NEXT: dead [[S_ADD_I32_9:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; FENCE-NEXT: ATOMIC_FENCE 5, 2
+ ; FENCE-NEXT: dead early-clobber %28:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; FENCE-NEXT: dead early-clobber %29:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; FENCE-NEXT: dead early-clobber %30:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; FENCE-NEXT: dead early-clobber %31:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; FENCE-NEXT: dead early-clobber %32:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; FENCE-NEXT: S_ENDPGM 0, implicit %10, implicit %11
+ ;
+ ; ALL-LABEL: name: test-block-carried-latency
+ ; ALL: bb.0:
+ ; ALL-NEXT: successors: %bb.1(0x80000000)
+ ; ALL-NEXT: liveins: $vgpr0, $sgpr0, $sgpr1
+ ; ALL-NEXT: {{ $}}
+ ; ALL-NEXT: [[DS_READ_B32_gfx9_:%[0-9]+]]:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 0, 0, implicit $exec
+ ; ALL-NEXT: [[DS_READ_B32_gfx9_1:%[0-9]+]]:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 4, 0, implicit $exec
+ ; ALL-NEXT: [[DS_READ_B32_gfx9_2:%[0-9]+]]:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 8, 0, implicit $exec
+ ; ALL-NEXT: [[DS_READ_B32_gfx9_3:%[0-9]+]]:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 12, 0, implicit $exec
+ ; ALL-NEXT: [[DEF:%[0-9]+]]:vreg_512_align2 = IMPLICIT_DEF
+ ; ALL-NEXT: [[DEF1:%[0-9]+]]:vreg_512_align2 = IMPLICIT_DEF
+ ; ALL-NEXT: [[DEF2:%[0-9]+]]:vreg_256_align2 = IMPLICIT_DEF
+ ; ALL-NEXT: [[DEF3:%[0-9]+]]:vreg_512_align2 = IMPLICIT_DEF
+ ; ALL-NEXT: [[DEF4:%[0-9]+]]:vreg_512_align2 = IMPLICIT_DEF
+ ; ALL-NEXT: [[DEF5:%[0-9]+]]:vreg_256_align2 = IMPLICIT_DEF
+ ; ALL-NEXT: S_BRANCH %bb.1
+ ; ALL-NEXT: {{ $}}
+ ; ALL-NEXT: bb.1:
+ ; ALL-NEXT: liveins: $vgpr0, $sgpr0, $sgpr1
+ ; ALL-NEXT: {{ $}}
+ ; ALL-NEXT: dead [[S_ADD_I32_:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; ALL-NEXT: dead [[S_ADD_I32_1:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; ALL-NEXT: dead [[S_ADD_I32_2:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; ALL-NEXT: dead [[S_ADD_I32_3:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; ALL-NEXT: dead [[V_ADD_U32_e64_:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; ALL-NEXT: dead [[S_ADD_I32_4:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; ALL-NEXT: dead [[V_ADD_U32_e64_1:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; ALL-NEXT: dead [[S_ADD_I32_5:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; ALL-NEXT: dead [[V_ADD_U32_e64_2:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; ALL-NEXT: dead [[S_ADD_I32_6:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; ALL-NEXT: dead [[V_ADD_U32_e64_3:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; ALL-NEXT: dead [[S_ADD_I32_7:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; ALL-NEXT: dead [[V_ADD_U32_e64_4:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; ALL-NEXT: dead [[S_ADD_I32_8:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; ALL-NEXT: dead [[V_ADD_U32_e64_5:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; ALL-NEXT: dead [[S_ADD_I32_9:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; ALL-NEXT: early-clobber %10:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF]], [[DEF1]], 0, [[DEF2]], [[DS_READ_B32_gfx9_]], [[DS_READ_B32_gfx9_1]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; ALL-NEXT: ATOMIC_FENCE 5, 2
+ ; ALL-NEXT: early-clobber %11:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; ALL-NEXT: dead early-clobber %28:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; ALL-NEXT: dead early-clobber %29:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; ALL-NEXT: dead early-clobber %30:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; ALL-NEXT: dead early-clobber %31:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; ALL-NEXT: dead early-clobber %32:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; ALL-NEXT: S_ENDPGM 0, implicit %10, implicit %11
+ bb.0:
+ successors: %bb.1
+ liveins: $vgpr0, $sgpr0, $sgpr1
+
+ %0:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 0, 0, implicit $exec
+ %1:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 4, 0, implicit $exec
+ %2:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 8, 0, implicit $exec
+ %3:vgpr_32_lo256 = DS_READ_B32_gfx9 $vgpr0, 12, 0, implicit $exec
+
+ %10:vreg_512_align2 = IMPLICIT_DEF
+ %11:vreg_512_align2 = IMPLICIT_DEF
+ %12:vreg_256_align2 = IMPLICIT_DEF
+ %15:vreg_512_align2 = IMPLICIT_DEF
+ %16:vreg_512_align2 = IMPLICIT_DEF
+ %17:vreg_256_align2 = IMPLICIT_DEF
+
+ S_BRANCH %bb.1
+
+ bb.1:
+ liveins: $vgpr0, $sgpr0, $sgpr1
+
+ ATOMIC_FENCE 5, 2
+
+ %20:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr %10, %11, 0, %12, %0, %1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ %21:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr %15, %16, 0, %17, %2, %3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+
+ %22:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ %23:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ %24:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ %25:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ %26:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ %27:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ %28:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ %29:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ %30:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ %31:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ %32:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ %33:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ %34:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ %35:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ %36:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ %37:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ %38:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr %15, %16, 0, %17, %2, %3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ %39:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr %15, %16, 0, %17, %2, %3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ %40:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr %15, %16, 0, %17, %2, %3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ %41:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr %15, %16, 0, %17, %2, %3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ %42:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr %15, %16, 0, %17, %2, %3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ S_ENDPGM 0, implicit %20, implicit %21
+...
>From 0677440179d0bf98af9e53b441bdd19dc58a83de Mon Sep 17 00:00:00 2001
From: Jeffrey Byrnes <Jeffrey.Byrnes at amd.com>
Date: Thu, 28 May 2026 15:32:38 -0700
Subject: [PATCH 09/10] Update test
Change-Id: I65d4e77a81c65f2f895cb295b74185a70b403de3
---
.../AMDGPU/coexec-block-carried-latency.mir | 28 ++--
llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll | 128 +++++++++---------
2 files changed, 76 insertions(+), 80 deletions(-)
diff --git a/llvm/test/CodeGen/AMDGPU/coexec-block-carried-latency.mir b/llvm/test/CodeGen/AMDGPU/coexec-block-carried-latency.mir
index 289a59fd1f032..4aa574ab54f82 100644
--- a/llvm/test/CodeGen/AMDGPU/coexec-block-carried-latency.mir
+++ b/llvm/test/CodeGen/AMDGPU/coexec-block-carried-latency.mir
@@ -38,21 +38,21 @@ body: |
; DEFAULT-NEXT: dead [[S_ADD_I32_1:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
; DEFAULT-NEXT: dead [[S_ADD_I32_2:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
; DEFAULT-NEXT: dead [[S_ADD_I32_3:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
- ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
; DEFAULT-NEXT: dead [[S_ADD_I32_4:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
; DEFAULT-NEXT: dead [[V_ADD_U32_e64_1:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; DEFAULT-NEXT: early-clobber %11:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
; DEFAULT-NEXT: dead [[S_ADD_I32_5:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
- ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_2:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
; DEFAULT-NEXT: dead [[S_ADD_I32_6:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
- ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_3:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_2:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
; DEFAULT-NEXT: dead [[S_ADD_I32_7:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
- ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_4:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
; DEFAULT-NEXT: dead [[S_ADD_I32_8:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
- ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_5:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
- ; DEFAULT-NEXT: early-clobber %11:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_3:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_4:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead early-clobber %28:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
; DEFAULT-NEXT: dead [[S_ADD_I32_9:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
; DEFAULT-NEXT: ATOMIC_FENCE 5, 2
- ; DEFAULT-NEXT: dead early-clobber %28:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; DEFAULT-NEXT: dead [[V_ADD_U32_e64_5:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
; DEFAULT-NEXT: dead early-clobber %29:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
; DEFAULT-NEXT: dead early-clobber %30:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
; DEFAULT-NEXT: dead early-clobber %31:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
@@ -84,21 +84,21 @@ body: |
; FENCE-NEXT: dead [[S_ADD_I32_1:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
; FENCE-NEXT: dead [[S_ADD_I32_2:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
; FENCE-NEXT: dead [[S_ADD_I32_3:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
- ; FENCE-NEXT: dead [[V_ADD_U32_e64_:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
; FENCE-NEXT: dead [[S_ADD_I32_4:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; FENCE-NEXT: dead [[V_ADD_U32_e64_:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
; FENCE-NEXT: dead [[V_ADD_U32_e64_1:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; FENCE-NEXT: early-clobber %11:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
; FENCE-NEXT: dead [[S_ADD_I32_5:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
- ; FENCE-NEXT: dead [[V_ADD_U32_e64_2:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
; FENCE-NEXT: dead [[S_ADD_I32_6:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
- ; FENCE-NEXT: dead [[V_ADD_U32_e64_3:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; FENCE-NEXT: dead [[V_ADD_U32_e64_2:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
; FENCE-NEXT: dead [[S_ADD_I32_7:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
- ; FENCE-NEXT: dead [[V_ADD_U32_e64_4:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
; FENCE-NEXT: dead [[S_ADD_I32_8:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
- ; FENCE-NEXT: dead [[V_ADD_U32_e64_5:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
- ; FENCE-NEXT: early-clobber %11:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
+ ; FENCE-NEXT: dead [[V_ADD_U32_e64_3:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; FENCE-NEXT: dead [[V_ADD_U32_e64_4:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
+ ; FENCE-NEXT: dead early-clobber %28:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
; FENCE-NEXT: dead [[S_ADD_I32_9:%[0-9]+]]:sgpr_32 = S_ADD_I32 $sgpr0, $sgpr1, implicit-def dead $scc
+ ; FENCE-NEXT: dead [[V_ADD_U32_e64_5:%[0-9]+]]:vgpr_32 = V_ADD_U32_e64 $vgpr0, $vgpr0, 0, implicit $exec
; FENCE-NEXT: ATOMIC_FENCE 5, 2
- ; FENCE-NEXT: dead early-clobber %28:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
; FENCE-NEXT: dead early-clobber %29:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
; FENCE-NEXT: dead early-clobber %30:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
; FENCE-NEXT: dead early-clobber %31:vreg_256_align2 = V_WMMA_SCALE_F32_16X16X128_F8F6F4_f8_f8_w32_threeaddr [[DEF3]], [[DEF4]], 0, [[DEF5]], [[DS_READ_B32_gfx9_2]], [[DS_READ_B32_gfx9_3]], 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, implicit $exec
diff --git a/llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll b/llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll
index 096ec69c00029..97a0d6880c6b6 100644
--- a/llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll
+++ b/llvm/test/CodeGen/AMDGPU/coexec-scheduler.ll
@@ -42,6 +42,9 @@ define amdgpu_kernel void @ds_wmma(ptr addrspace(3) %base, ptr addrspace(1) %out
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_mov_b32_e32 v88, s2
+; COEXEC-NEXT: s_add_co_i32 s2, s2, s1
+; COEXEC-NEXT: s_and_b32 s3, s0, exec_lo
+; COEXEC-NEXT: s_cselect_b32 s3, 1, 0
; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v88 offset:192
; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v88 offset:128
; COEXEC-NEXT: ds_load_tr16_b128 v[40:43], v88
@@ -58,9 +61,6 @@ define amdgpu_kernel void @ds_wmma(ptr addrspace(3) %base, ptr addrspace(1) %out
; COEXEC-NEXT: ds_load_tr16_b128 v[80:83], v88 offset:896
; COEXEC-NEXT: ds_load_tr16_b128 v[92:95], v88 offset:832
; COEXEC-NEXT: ds_load_tr16_b128 v[88:91], v88 offset:768
-; COEXEC-NEXT: s_add_co_i32 s2, s2, s1
-; COEXEC-NEXT: s_and_b32 s3, s0, exec_lo
-; COEXEC-NEXT: s_cselect_b32 s3, 1, 0
; COEXEC-NEXT: s_cmp_lg_u32 s3, 1
; COEXEC-NEXT: s_wait_dscnt 0xc
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[32:39], v[24:31]
@@ -286,73 +286,73 @@ define amdgpu_kernel void @ds_wmma_permute(ptr addrspace(3) %base, ptr addrspace
; COEXEC-NEXT: .LBB1_1: ; %loop
; COEXEC-NEXT: ; =>This Inner Loop Header: Depth=1
; COEXEC-NEXT: s_add_co_i32 s7, s0, s6
+; COEXEC-NEXT: s_add_co_i32 s8, s1, s6
+; COEXEC-NEXT: s_add_co_i32 s6, s6, s3
+; COEXEC-NEXT: s_and_b32 s9, s2, exec_lo
+; COEXEC-NEXT: s_cselect_b32 s9, 1, 0
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
-; COEXEC-NEXT: v_mov_b32_e32 v92, s7
-; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v92 offset:64
-; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v92
-; COEXEC-NEXT: ds_load_tr16_b128 v[40:43], v92 offset:128
-; COEXEC-NEXT: ds_load_tr16_b128 v[44:47], v92 offset:192
-; COEXEC-NEXT: ds_load_tr16_b128 v[48:51], v92 offset:256
-; COEXEC-NEXT: ds_load_tr16_b128 v[52:55], v92 offset:320
-; COEXEC-NEXT: ds_load_tr16_b128 v[56:59], v92 offset:384
-; COEXEC-NEXT: ds_load_tr16_b128 v[60:63], v92 offset:448
-; COEXEC-NEXT: ds_load_tr16_b128 v[64:67], v92 offset:512
-; COEXEC-NEXT: ds_load_tr16_b128 v[68:71], v92 offset:576
-; COEXEC-NEXT: ds_load_tr16_b128 v[72:75], v92 offset:640
-; COEXEC-NEXT: ds_load_tr16_b128 v[76:79], v92 offset:704
-; COEXEC-NEXT: ds_load_tr16_b128 v[80:83], v92 offset:768
-; COEXEC-NEXT: ds_load_tr16_b128 v[84:87], v92 offset:832
-; COEXEC-NEXT: ds_load_tr16_b128 v[88:91], v92 offset:896
-; COEXEC-NEXT: ds_load_tr16_b128 v[92:95], v92 offset:960
-; COEXEC-NEXT: s_add_co_i32 s7, s1, s6
-; COEXEC-NEXT: s_add_co_i32 s6, s6, s3
-; COEXEC-NEXT: s_and_b32 s8, s2, exec_lo
-; COEXEC-NEXT: s_cselect_b32 s8, 1, 0
-; COEXEC-NEXT: v_mov_b32_e32 v156, s7
-; COEXEC-NEXT: s_cmp_lg_u32 s8, 1
-; COEXEC-NEXT: ds_load_tr16_b128 v[96:99], v156
-; COEXEC-NEXT: ds_load_tr16_b128 v[100:103], v156 offset:64
-; COEXEC-NEXT: ds_load_tr16_b128 v[108:111], v156 offset:320
-; COEXEC-NEXT: ds_load_tr16_b128 v[104:107], v156 offset:256
-; COEXEC-NEXT: ds_load_tr16_b128 v[116:119], v156 offset:576
-; COEXEC-NEXT: ds_load_tr16_b128 v[112:115], v156 offset:512
-; COEXEC-NEXT: ds_load_tr16_b128 v[124:127], v156 offset:832
-; COEXEC-NEXT: ds_load_tr16_b128 v[120:123], v156 offset:768
+; COEXEC-NEXT: v_mov_b32_e32 v124, s7
+; COEXEC-NEXT: s_cmp_lg_u32 s9, 1
+; COEXEC-NEXT: v_mov_b32_e32 v156, s8
+; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v124
+; COEXEC-NEXT: ds_load_tr16_b128 v[36:39], v124 offset:64
+; COEXEC-NEXT: ds_load_tr16_b128 v[40:43], v156
+; COEXEC-NEXT: ds_load_tr16_b128 v[44:47], v156 offset:64
+; COEXEC-NEXT: ds_load_tr16_b128 v[52:55], v124 offset:320
+; COEXEC-NEXT: ds_load_tr16_b128 v[48:51], v124 offset:256
+; COEXEC-NEXT: ds_load_tr16_b128 v[56:59], v156 offset:256
+; COEXEC-NEXT: ds_load_tr16_b128 v[60:63], v156 offset:320
+; COEXEC-NEXT: ds_load_tr16_b128 v[68:71], v124 offset:576
+; COEXEC-NEXT: ds_load_tr16_b128 v[64:67], v124 offset:512
+; COEXEC-NEXT: ds_load_tr16_b128 v[76:79], v156 offset:576
+; COEXEC-NEXT: ds_load_tr16_b128 v[72:75], v156 offset:512
+; COEXEC-NEXT: ds_load_tr16_b128 v[84:87], v124 offset:832
+; COEXEC-NEXT: ds_load_tr16_b128 v[80:83], v124 offset:768
+; COEXEC-NEXT: ds_load_tr16_b128 v[92:95], v156 offset:832
+; COEXEC-NEXT: ds_load_tr16_b128 v[88:91], v156 offset:768
+; COEXEC-NEXT: ds_load_tr16_b128 v[96:99], v124 offset:128
+; COEXEC-NEXT: ds_load_tr16_b128 v[100:103], v124 offset:192
+; COEXEC-NEXT: ds_load_tr16_b128 v[104:107], v124 offset:384
+; COEXEC-NEXT: s_wait_dscnt 0xf
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[40:47], v[24:31]
+; COEXEC-NEXT: ds_load_tr16_b128 v[108:111], v124 offset:448
+; COEXEC-NEXT: ds_load_tr16_b128 v[112:115], v124 offset:640
+; COEXEC-NEXT: ds_load_tr16_b128 v[116:119], v124 offset:704
+; COEXEC-NEXT: ds_load_tr16_b128 v[120:123], v124 offset:896
+; COEXEC-NEXT: ds_load_tr16_b128 v[124:127], v124 offset:960
; COEXEC-NEXT: ds_load_tr16_b128 v[128:131], v156 offset:128
; COEXEC-NEXT: ds_load_tr16_b128 v[132:135], v156 offset:192
+; COEXEC-NEXT: s_wait_dscnt 0x12
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[56:63], v[16:23]
; COEXEC-NEXT: ds_load_tr16_b128 v[136:139], v156 offset:384
; COEXEC-NEXT: ds_load_tr16_b128 v[140:143], v156 offset:448
; COEXEC-NEXT: ds_load_tr16_b128 v[144:147], v156 offset:640
; COEXEC-NEXT: ds_load_tr16_b128 v[148:151], v156 offset:704
; COEXEC-NEXT: ds_load_tr16_b128 v[152:155], v156 offset:896
; COEXEC-NEXT: ds_load_tr16_b128 v[156:159], v156 offset:960
-; COEXEC-NEXT: s_wait_dscnt 0xe
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[96:103], v[24:31]
-; COEXEC-NEXT: s_wait_dscnt 0xc
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[104:111], v[16:23]
-; COEXEC-NEXT: s_wait_dscnt 0xa
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[112:119], v[8:15]
-; COEXEC-NEXT: s_wait_dscnt 0x8
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[80:87], v[120:127], v[0:7]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[96:103], v[24:31]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[104:111], v[16:23]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[112:119], v[8:15]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[80:87], v[120:127], v[0:7]
+; COEXEC-NEXT: s_wait_dscnt 0x14
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[72:79], v[8:15]
+; COEXEC-NEXT: s_wait_dscnt 0x10
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[80:87], v[88:95], v[0:7]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[40:47], v[24:31]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[56:63], v[16:23]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[72:79], v[8:15]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[80:87], v[88:95], v[0:7]
; COEXEC-NEXT: s_wait_dscnt 0x6
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[128:135], v[24:31]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[96:103], v[128:135], v[24:31]
; COEXEC-NEXT: s_wait_dscnt 0x4
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[56:63], v[136:143], v[16:23]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[104:111], v[136:143], v[16:23]
; COEXEC-NEXT: s_wait_dscnt 0x2
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[72:79], v[144:151], v[8:15]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[112:119], v[144:151], v[8:15]
; COEXEC-NEXT: s_wait_dscnt 0x0
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[88:95], v[152:159], v[0:7]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[40:47], v[128:135], v[24:31]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[56:63], v[136:143], v[16:23]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[72:79], v[144:151], v[8:15]
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[88:95], v[152:159], v[0:7]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[120:127], v[152:159], v[0:7]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[96:103], v[128:135], v[24:31]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[104:111], v[136:143], v[16:23]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[112:119], v[144:151], v[8:15]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[120:127], v[152:159], v[0:7]
; COEXEC-NEXT: s_cbranch_scc1 .LBB1_1
; COEXEC-NEXT: ; %bb.2: ; %end
; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
@@ -657,7 +657,8 @@ define amdgpu_kernel void @ds_wmma_block_carried(ptr addrspace(3) %base, ptr add
; COEXEC-NEXT: s_wait_dscnt 0x0
; COEXEC-NEXT: .LBB2_1: ; %loop
; COEXEC-NEXT: ; =>This Inner Loop Header: Depth=1
-; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: s_wait_dscnt 0x4
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[40:47], v[24:31]
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
@@ -669,13 +670,10 @@ define amdgpu_kernel void @ds_wmma_block_carried(ptr addrspace(3) %base, ptr add
; COEXEC-NEXT: s_add_co_i32 s0, s0, s1
; COEXEC-NEXT: s_and_b32 s3, s2, exec_lo
; COEXEC-NEXT: s_cselect_b32 s3, 1, 0
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[64:71], v[8:15]
; COEXEC-NEXT: s_cmp_lg_u32 s3, 1
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[64:71], v[8:15]
-; COEXEC-NEXT: s_wait_dscnt 0x4
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[40:47], v[24:31]
; COEXEC-NEXT: s_wait_dscnt 0x0
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[16:23], v[48:55], v[56:63], v[16:23]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[64:71], v[8:15]
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[64:71], v[64:71], v[0:7]
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[24:31], v[32:39], v[40:47], v[24:31]
; COEXEC-NEXT: ds_load_tr16_b128 v[32:35], v72
@@ -687,13 +685,13 @@ define amdgpu_kernel void @ds_wmma_block_carried(ptr addrspace(3) %base, ptr add
; COEXEC-NEXT: ds_load_tr16_b128 v[52:55], v72 offset:320
; COEXEC-NEXT: ds_load_tr16_b128 v[56:59], v72 offset:384
; COEXEC-NEXT: ds_load_tr16_b128 v[60:63], v72 offset:448
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[8:15], v[64:71], v[64:71], v[8:15]
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[0:7], v[64:71], v[64:71], v[0:7]
; COEXEC-NEXT: s_cbranch_scc1 .LBB2_1
; COEXEC-NEXT: ; %bb.2: ; %end
; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
; COEXEC-NEXT: s_wait_dscnt 0x7
; COEXEC-NEXT: v_nop
-; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_mov_b32_e32 v32, 0
; COEXEC-NEXT: s_wait_kmcnt 0x0
; COEXEC-NEXT: s_clause 0x7
@@ -970,7 +968,8 @@ define amdgpu_kernel void @ds_wmma_loop_carried(ptr addrspace(3) %base, ptr addr
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[88:95], v[64:71], v[32:39]
; COEXEC-NEXT: .LBB3_1: ; %loop
; COEXEC-NEXT: ; =>This Inner Loop Header: Depth=1
-; COEXEC-NEXT: v_nop
+; COEXEC-NEXT: s_wait_dscnt 0x4
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[0:7], v[16:23], v[56:63]
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_nop
@@ -982,13 +981,10 @@ define amdgpu_kernel void @ds_wmma_loop_carried(ptr addrspace(3) %base, ptr addr
; COEXEC-NEXT: s_add_co_i32 s0, s0, s1
; COEXEC-NEXT: s_and_b32 s3, s2, exec_lo
; COEXEC-NEXT: s_cselect_b32 s3, 1, 0
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[40:47], v[64:71], v[64:71], v[40:47]
; COEXEC-NEXT: s_cmp_lg_u32 s3, 1
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[40:47], v[64:71], v[64:71], v[40:47]
-; COEXEC-NEXT: s_wait_dscnt 0x4
-; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[0:7], v[16:23], v[56:63]
; COEXEC-NEXT: s_wait_dscnt 0x0
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[48:55], v[8:15], v[24:31], v[48:55]
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[40:47], v[64:71], v[64:71], v[40:47]
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[64:71], v[64:71], v[32:39]
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[56:63], v[0:7], v[16:23], v[56:63]
; COEXEC-NEXT: ds_load_tr16_b128 v[0:3], v72
@@ -1000,13 +996,13 @@ define amdgpu_kernel void @ds_wmma_loop_carried(ptr addrspace(3) %base, ptr addr
; COEXEC-NEXT: ds_load_tr16_b128 v[12:15], v72 offset:320
; COEXEC-NEXT: ds_load_tr16_b128 v[24:27], v72 offset:384
; COEXEC-NEXT: ds_load_tr16_b128 v[28:31], v72 offset:448
+; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[40:47], v[64:71], v[64:71], v[40:47]
; COEXEC-NEXT: v_wmma_f32_16x16x32_f16 v[32:39], v[64:71], v[64:71], v[32:39]
; COEXEC-NEXT: s_cbranch_scc1 .LBB3_1
; COEXEC-NEXT: ; %bb.2: ; %end
; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x8 nv
; COEXEC-NEXT: s_wait_dscnt 0x7
; COEXEC-NEXT: v_nop
-; COEXEC-NEXT: v_nop
; COEXEC-NEXT: v_mov_b32_e32 v0, 0
; COEXEC-NEXT: s_wait_kmcnt 0x0
; COEXEC-NEXT: s_clause 0x7
>From 09fc36fd58518bc4a7a90cdb31b74ae33edcccb8 Mon Sep 17 00:00:00 2001
From: Austin Kerbow <Austin.Kerbow at amd.com>
Date: Fri, 11 Sep 2026 13:07:05 -0500
Subject: [PATCH 10/10] Address comments.
---
.../AMDGPU/AMDGPUCoExecSchedStrategy.cpp | 24 ++++++++++---------
.../Target/AMDGPU/AMDGPUCoExecSchedStrategy.h | 10 ++++----
.../AMDGPU/coexec-mfma-interleave-gfx950.ll | 10 ++++----
.../CodeGen/AMDGPU/coexec-sched-ds-fifo.mir | 2 +-
.../llvm.amdgcn.sched.group.barrier.gfx12.ll | 5 ++--
5 files changed, 26 insertions(+), 25 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
index 2e0a85ccb5fd5..6a098f36c43fa 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.cpp
@@ -26,10 +26,11 @@ enum class CarriedLatency { Off, Fence, All };
static cl::opt<CarriedLatency> BlockCarriedLatency(
"amdgpu-block-carried-latency", cl::Hidden, cl::init(CarriedLatency::Off),
- cl::desc("Prioritize HardwareUnits with non-zero exposed cycles in the "
- "coexec scheduler's critical-resource sort."),
+ cl::desc("Estimate block-carried latency and include it in the effective "
+ "candidate stall cost."),
cl::values(
- clEnumValN(CarriedLatency::Off, "off", "Disabled- do not pad latency."),
+ clEnumValN(CarriedLatency::Off, "off",
+ "Disabled - do not pad latency."),
clEnumValN(CarriedLatency::Fence, "fence",
"Only pad latency for memory fence (e.g. those surrounding "
"barrier_signal/wait)."),
@@ -562,7 +563,7 @@ CandidateHeuristics::getHWUIFromFlavor(InstructionFlavor Flavor) {
unsigned CandidateHeuristics::getMaxBlockingCycles(const MCSchedClassDesc *SC,
const MachineInstr *MI) {
- // Loads and stores are not pipelined
+ // Loads and stores are not pipelined.
if (MI->mayLoadOrStore())
return SchedModel->computeInstrLatency(MI, false);
@@ -630,11 +631,12 @@ unsigned CandidateHeuristics::getCarriedLatency(SUnit *SU) {
return 0;
MachineInstr *MI = SU->getInstr();
+ unsigned CarriedLatency = 0;
const InstructionFlavor Flavor = classifyFlavor(*MI, *SII);
if (Flavor == InstructionFlavor::Fence) {
MachineBasicBlock *MBB = MI->getParent();
- // Check if we have DS instruciton after a fence in any of the predecessor
- // blocks, if so, this fence instruction has carried latency.
+ // Scan each direct predecessor back to its nearest Fence or block start for
+ // DS instructions.
for (auto PredMBB : MBB->predecessors()) {
auto I = PredMBB->rbegin();
auto E = PredMBB->rend();
@@ -645,15 +647,14 @@ unsigned CandidateHeuristics::getCarriedLatency(SUnit *SU) {
// Found carried latency.
if (ItFlavor == InstructionFlavor::DS)
- return getHWUICyclesForMI(&*I);
+ CarriedLatency = std::max(CarriedLatency, getHWUICyclesForMI(&*I));
}
}
}
if (BlockCarriedLatency == CarriedLatency::Fence)
- return 0;
+ return CarriedLatency;
- unsigned CarriedLatency = 0;
for (MachineOperand &Op : MI->all_uses()) {
auto Reg = Op.getReg();
if (!Reg.isVirtual())
@@ -668,7 +669,7 @@ unsigned CandidateHeuristics::getCarriedLatency(SUnit *SU) {
unsigned Latency = getHWUICyclesForMI(&Def);
- // Load is carried across block
+ // Load is carried across block.
if (Def.getParent() != MI->getParent()) {
bool FoundUseInDefBlock = false;
for (MachineInstr &Use : DAG->MRI.use_nodbg_instructions(Reg)) {
@@ -692,7 +693,7 @@ unsigned CandidateHeuristics::getCarriedLatency(SUnit *SU) {
}
assert(Def.getParent() == MI->getParent());
- // Load is in the same block
+ // Load is in the same block.
SlotIndex LoadIdx = DAG->getLIS()->getInstructionIndex(Def);
SlotIndex UseIdx = DAG->getLIS()->getInstructionIndex(*MI);
// The load occurs after this use -- the latency is carried across loop
@@ -705,6 +706,7 @@ unsigned CandidateHeuristics::getCarriedLatency(SUnit *SU) {
}
void CandidateHeuristics::collectRegionSummary() {
+ CarriedLatencies.clear();
if (!SchedModel || !SchedModel->hasInstrSchedModel())
return;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.h b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.h
index 6e372f52d2069..9032df3724ab1 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCoExecSchedStrategy.h
@@ -210,15 +210,15 @@ class CandidateHeuristics {
void collectRegionSummary();
/// \returns the maximum blocking cycles according to the SchedModel for a
- /// given MCSchedClassDesc \p SC
+ /// given MCSchedClassDesc \p SC.
unsigned getMaxBlockingCycles(const MCSchedClassDesc *SC,
const MachineInstr *MI);
/// Compute the blocking cycles for the appropriate HardwareUnit given an \p
- /// SU
+ /// SU.
unsigned getHWUICyclesForSU(SUnit *SU);
/// Compute the blocking cycles for the appropriate HardwareUnit given an \p
- /// MI
+ /// MI.
unsigned getHWUICyclesForMI(MachineInstr *MI);
/// Estimate the block carried latency from loads for a given \p SU. This is
@@ -248,8 +248,8 @@ class CandidateHeuristics {
unsigned getStructuralStallCycles(SchedBoundary &Zone, SUnit *SU);
- bool tryEffectiveStall(GenericSchedulerBase::SchedCandidate &TryCand,
- GenericSchedulerBase::SchedCandidate &Cand,
+ bool tryEffectiveStall(GenericSchedulerBase::SchedCandidate &Cand,
+ GenericSchedulerBase::SchedCandidate &TryCand,
SchedBoundary &Zone);
/// Check for critical resource consumption. Prefer the candidate that uses
diff --git a/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll b/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll
index c290874a5d89b..3ebfd4d1d5a24 100644
--- a/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll
+++ b/llvm/test/CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll
@@ -11,12 +11,11 @@ define amdgpu_kernel void @mfma_16x16_interleave(
; CHECK-LABEL: mfma_16x16_interleave:
; CHECK: ; %bb.0:
; CHECK-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x0
-; CHECK-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x8
-; CHECK-NEXT: v_and_b32_e32 v2, 0x3ff, v0
+; CHECK-NEXT: v_and_b32_e32 v14, 0x3ff, v0
; CHECK-NEXT: v_mov_b32_e32 v1, 0
-; CHECK-NEXT: v_lshlrev_b32_e32 v0, 4, v2
-; CHECK-NEXT: v_mul_hi_i32_i24_e32 v15, -12, v2
-; CHECK-NEXT: v_mul_i32_i24_e32 v14, -12, v2
+; CHECK-NEXT: v_lshlrev_b32_e32 v0, 4, v14
+; CHECK-NEXT: v_mul_hi_i32_i24_e32 v15, -12, v14
+; CHECK-NEXT: v_mul_i32_i24_e32 v14, -12, v14
; CHECK-NEXT: s_waitcnt lgkmcnt(0)
; CHECK-NEXT: v_lshl_add_u64 v[16:17], s[0:1], 0, v[0:1]
; CHECK-NEXT: global_load_dwordx4 v[2:5], v0, s[0:1]
@@ -25,6 +24,7 @@ define amdgpu_kernel void @mfma_16x16_interleave(
; CHECK-NEXT: v_lshl_add_u64 v[18:19], v[16:17], 0, v[14:15]
; CHECK-NEXT: global_load_dwordx2 v[30:31], v[18:19], off offset:16
; CHECK-NEXT: global_load_dwordx4 v[14:17], v[18:19], off
+; CHECK-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x8
; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; CHECK-NEXT: ; sched_barrier mask(0x00000000)
; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[18:21], v[2:5], v[6:9], v[2:5]
diff --git a/llvm/test/CodeGen/AMDGPU/coexec-sched-ds-fifo.mir b/llvm/test/CodeGen/AMDGPU/coexec-sched-ds-fifo.mir
index 05a08205ea767..7348084dedc77 100644
--- a/llvm/test/CodeGen/AMDGPU/coexec-sched-ds-fifo.mir
+++ b/llvm/test/CodeGen/AMDGPU/coexec-sched-ds-fifo.mir
@@ -2,7 +2,7 @@
# RUN: llc -mtriple=amdgpu12.50 -run-pass=machine-scheduler -amdgpu-sched-strategy=coexec -debug-only=machine-scheduler %s -filetype=null 2>&1 | FileCheck --check-prefix=DEBUG %s
# RUN: llc -mtriple=amdgpu12.50 -run-pass=machine-scheduler -amdgpu-sched-strategy=coexec %s -o - | FileCheck --check-prefix=ORDER %s
-# DEBUG: Effective stalls: try=4 (ready=0, struct=0, lat=0, buffer=4) cand=0 (ready=0, struct=0, lat=0, buffer=0)
+# DEBUG: Effective stalls: try=4 (ready=0, struct=0, lat=0, carried=0, buffer=4) cand=0 (ready=0, struct=0, lat=0, carried=0, buffer=0)
# ORDER-LABEL: name: ds_fifo_stall
# ORDER: body: |
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.gfx12.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.gfx12.ll
index a28b44e9fa38d..38106eb2017a2 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.gfx12.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.sched.group.barrier.gfx12.ll
@@ -108,9 +108,8 @@ define amdgpu_kernel void @test_sched_group_barrier_pipeline_SWMMAC_cluster(ptr
; COEXEC-NEXT: s_mov_b64 s[64:65], 0
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
-; COEXEC-NEXT: v_mov_b32_e32 v48, 0
; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv
-; COEXEC-NEXT: v_lshlrev_b32_e32 v0, 4, v0
+; COEXEC-NEXT: v_dual_mov_b32 v48, 0 :: v_dual_lshlrev_b32 v0, 4, v0
; COEXEC-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
; COEXEC-NEXT: v_and_b32_e32 v0, 0x3ff0, v0
; COEXEC-NEXT: s_wait_kmcnt 0x0
@@ -329,8 +328,8 @@ define amdgpu_kernel void @test_sched_group_barrier_pipeline_SWMMAC_interleaved(
; COEXEC-NEXT: s_mov_b64 s[64:65], 0
; COEXEC-NEXT: v_nop
; COEXEC-NEXT: global_prefetch_b8 v0, s[64:65] scope:SCOPE_SE
-; COEXEC-NEXT: v_mov_b32_e32 v16, 0
; COEXEC-NEXT: s_load_b64 s[0:1], s[4:5], 0x24 nv
+; COEXEC-NEXT: v_mov_b32_e32 v16, 0
; COEXEC-NEXT: v_and_b32_e32 v0, 0x3ff, v0
; COEXEC-NEXT: s_wait_kmcnt 0x0
; COEXEC-NEXT: v_mov_b32_e32 v17, s1
More information about the llvm-commits
mailing list