[llvm] [AMDGPU] Serialize disjoint MFMA chains to hide DS_READ latency (PR #170242)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Dec 1 20:36:39 PST 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Anshil Gandhi (gandhi56)
<details>
<summary>Changes</summary>
This patch identifies disjoint chains of dependent MFMA instructions (with length >= 2) and stitches them together into a single execution sequence by adding artificial dependencies from the tail of one chain to the head of the next.
Currently, the scheduler may schedule disjoint MFMA chains too early or interleave them, which can expose high latencies from their associated DS_READ operands. By strictly serializing these MFMA chains, we force subsequent chains to execute later. This artificial delay increases the distance between the DS_READ issuance and the consuming MFMA instruction, effectively hiding the load latency.
---
Patch is 158.84 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/170242.diff
5 Files Affected:
- (modified) llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp (+96-1)
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir (+757-750)
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.small.mir (+369-366)
- (added) llvm/test/CodeGen/AMDGPU/misched-ds-mfma-order-false-deps.mir (+118)
- (modified) llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll (+40-43)
``````````diff
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
index 85addb13aef8d..7a4f0ccf6c88a 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
@@ -21,6 +21,7 @@
#include "SIMachineFunctionInfo.h"
#include "llvm/ADT/BitmaskEnum.h"
#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/STLExtras.h"
#include "llvm/CodeGen/MachineScheduler.h"
#include "llvm/CodeGen/TargetOpcodes.h"
@@ -60,6 +61,10 @@ static cl::opt<bool> UseCostHeur(
"Experimentally, results are mixed, so this should be set on a "
"case-by-case basis."));
+static cl::opt<bool> DisableMfmaChainOrderingDeps(
+ "amdgpu-disable-mfma-chain-order-deps", cl::init(false), cl::Hidden,
+ cl::desc("Enable artificial false dependencies between MFMA chains"));
+
// Components of the mask that determines which instruction types may be may be
// classified into a SchedGroup.
enum class SchedGroupMask {
@@ -2342,6 +2347,10 @@ class IGroupLPDAGMutation : public ScheduleDAGMutation {
// Add DAG edges that enforce SCHED_BARRIER ordering.
void addSchedBarrierEdges(SUnit &SU);
+ // Add artificial false-dependencies between MFMA consumers of adjacent
+ // DS_READ_B128 streams to enforce MFMA(newer) -> MFMA(older-last) ordering.
+ void addMfmaFalseDeps();
+
// Use a SCHED_BARRIER's mask to identify instruction SchedGroups that should
// not be reordered accross the SCHED_BARRIER. This is used for the base
// SCHED_BARRIER, and not SCHED_GROUP_BARRIER. The difference is that
@@ -2585,6 +2594,9 @@ void IGroupLPDAGMutation::apply(ScheduleDAGInstrs *DAGInstrs) {
}
}
+ if (!DisableMfmaChainOrderingDeps && ST.hasMAIInsts())
+ addMfmaFalseDeps();
+
if (FoundSB || (FoundIGLP && ShouldApplyIGLP)) {
PipelineSolver PS(SyncedSchedGroups, SyncedInstrs, DAG, IsBottomUp);
// PipelineSolver performs the mutation by adding the edges it
@@ -2681,12 +2693,95 @@ bool IGroupLPDAGMutation::initIGLPOpt(SUnit &SU) {
} // namespace
+void IGroupLPDAGMutation::addMfmaFalseDeps() {
+ DenseMap<SUnit *, SUnit *> MFMAAncestor;
+ SmallVector<SUnit *, 10> MFMAChainLeaders;
+ DenseMap<SUnit *, SUnit *> MFMAChainNext;
+ for (auto &SU : DAG->SUnits) {
+ if (!TII->isMFMAorWMMA(*SU.getInstr()))
+ continue;
+
+ if (MFMAAncestor.contains(&SU))
+ continue;
+
+ SUnit *CurrMFMA = &SU;
+ MFMAAncestor[CurrMFMA] = CurrMFMA;
+ MFMAChainLeaders.push_back(&SU);
+ while (!CurrMFMA->Succs.empty()) {
+ // Count the number of successor MFMA/WMMA instructions of
+ // the current MFMA instruction.
+ SUnit *NextMFMA = nullptr;
+ unsigned MFMADataDepSuccCount = 0;
+ for (const auto &Succ : CurrMFMA->Succs) {
+ SUnit *SuccSU = Succ.getSUnit();
+ if (!SuccSU->isInstr() || !TII->isMFMAorWMMA(*SuccSU->getInstr()))
+ continue;
+
+ // Check if the successor is MFMA/WMMA and the edge is a data dependency
+ if (Succ.getKind() == SDep::Data) {
+ NextMFMA = SuccSU;
+ MFMADataDepSuccCount++;
+ }
+ }
+
+ // If the current MFMA instruction has more than one successor MFMA/WMMA instruction,
+ // we need to break the chain.
+ if (MFMADataDepSuccCount != 1) {
+ MFMAChainNext[CurrMFMA] = nullptr;
+ break;
+ }
+
+ // Add the current MFMA instruction to the MFMAAncestor map.
+ MFMAAncestor[CurrMFMA] = &SU;
+ MFMAChainNext[CurrMFMA] = NextMFMA;
+ CurrMFMA = NextMFMA;
+ }
+ }
+
+ // Compute the tail and length of each chain in a single loop.
+ auto GetTailAndLength = [&](SUnit *Leader) -> std::pair<SUnit *, unsigned> {
+ unsigned Length = 1;
+ SUnit *Curr = Leader;
+ while (MFMAChainNext.count(Curr)) {
+ if (!MFMAChainNext[Curr])
+ break;
+ Curr = MFMAChainNext[Curr];
+ ++Length;
+ }
+ return {Curr, Length};
+ };
+
+ // Assert that all MFMA chains are ordered by NodeNum
+ // Add artificial false dependencies between MFMA chains if two given
+ // chains are at least 2 SUs long.
+ // Iterate over all pairs of contiguous MFMA chains and add artificial edges if chains are at least 2 SUs long.
+ for (size_t I = 0; I + 1 < MFMAChainLeaders.size(); ++I) {
+ SUnit *ChainLeaderA = MFMAChainLeaders[I];
+ SUnit *ChainLeaderB = MFMAChainLeaders[I + 1];
+
+ auto [TailA, LengthA] = GetTailAndLength(ChainLeaderA);
+ auto [TailB, LengthB] = GetTailAndLength(ChainLeaderB);
+
+ // Only add if both chains are at least two SUs long.
+ if (LengthA >= 2 && LengthB >= 2) {
+ // Add an artificial dependency edge from the tail of chain A to the
+ // leader of chain B.
+ LLVM_DEBUG(dbgs() << "Adding artificial dependency edge from " << TailA->NodeNum
+ << " to " << ChainLeaderB->NodeNum << "\n");
+ DAG->addEdge(ChainLeaderB, SDep(TailA, SDep::Artificial));
+ }
+ }
+}
+
+namespace llvm {
+
/// \p Phase specifes whether or not this is a reentry into the
/// IGroupLPDAGMutation. Since there may be multiple scheduling passes on the
/// same scheduling region (e.g. pre and post-RA scheduling / multiple
/// scheduling "phases"), we can reenter this mutation framework more than once
/// for a given region.
std::unique_ptr<ScheduleDAGMutation>
-llvm::createIGroupLPDAGMutation(AMDGPU::SchedulingPhase Phase) {
+createIGroupLPDAGMutation(AMDGPU::SchedulingPhase Phase) {
return std::make_unique<IGroupLPDAGMutation>(Phase);
}
+}
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir
index 689d1472d6010..a2a00d107a7bc 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir
@@ -6,143 +6,143 @@
define amdgpu_kernel void @largeInterleave() #0 { ret void }
; GCN-LABEL: largeInterleave:
; GCN: ; %bb.0:
- ; GCN-NEXT: ; implicit-def: $vgpr16
- ; GCN-NEXT: ; implicit-def: $vgpr25
+ ; GCN-NEXT: ; implicit-def: $vgpr0
; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
- ; GCN-NEXT: v_readfirstlane_b32 s17, v16
+ ; GCN-NEXT: v_readfirstlane_b32 s17, v0
+ ; GCN-NEXT: ; implicit-def: $vgpr2
+ ; GCN-NEXT: s_lshl_b32 s18, s17, 7
+ ; GCN-NEXT: ; implicit-def: $vgpr4
+ ; GCN-NEXT: v_add_lshl_u32 v231, v2, s18, 1
+ ; GCN-NEXT: v_add_u32_e32 v2, s17, v4
+ ; GCN-NEXT: v_and_b32_e32 v2, 0x1fffffff, v2
+ ; GCN-NEXT: ; implicit-def: $sgpr16
+ ; GCN-NEXT: v_mul_lo_u32 v2, v2, s16
+ ; GCN-NEXT: ; implicit-def: $vgpr5
+ ; GCN-NEXT: ; implicit-def: $vgpr9
+ ; GCN-NEXT: v_add_lshl_u32 v172, v5, v2, 1
+ ; GCN-NEXT: v_lshl_add_u32 v2, s17, 4, v9
; GCN-NEXT: ; implicit-def: $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15
- ; GCN-NEXT: ; implicit-def: $vgpr17
+ ; GCN-NEXT: v_mul_lo_u32 v2, v2, s6
+ ; GCN-NEXT: ; implicit-def: $vgpr1
+ ; GCN-NEXT: v_add_lshl_u32 v224, v2, v1, 1
; GCN-NEXT: ; implicit-def: $sgpr15
+ ; GCN-NEXT: v_add_u32_e32 v1, s15, v224
; GCN-NEXT: ; implicit-def: $sgpr8_sgpr9_sgpr10_sgpr11
- ; GCN-NEXT: s_lshl_b32 s18, s17, 7
- ; GCN-NEXT: ; implicit-def: $vgpr18
- ; GCN-NEXT: v_add_lshl_u32 v230, v18, s18, 1
- ; GCN-NEXT: v_lshl_add_u32 v25, s17, 4, v25
- ; GCN-NEXT: v_mul_lo_u32 v25, v25, s6
- ; GCN-NEXT: v_add_lshl_u32 v226, v25, v17, 1
- ; GCN-NEXT: v_add_u32_e32 v17, s15, v226
- ; GCN-NEXT: buffer_load_dwordx4 v[64:67], v226, s[8:11], 0 offen sc0 sc1
+ ; GCN-NEXT: buffer_load_dwordx4 v[64:67], v224, s[8:11], 0 offen sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: buffer_load_dwordx4 v[68:71], v17, s[8:11], 0 offen sc0 sc1
+ ; GCN-NEXT: buffer_load_dwordx4 v[68:71], v1, s[8:11], 0 offen sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_add_u32_e32 v72, 64, v17
- ; GCN-NEXT: ; implicit-def: $vgpr213
+ ; GCN-NEXT: v_add_u32_e32 v72, 64, v1
+ ; GCN-NEXT: ; implicit-def: $vgpr230
; GCN-NEXT: ; implicit-def: $vgpr152_vgpr153_vgpr154_vgpr155
- ; GCN-NEXT: ; implicit-def: $vgpr246
- ; GCN-NEXT: v_add_u32_e32 v188, 0x80, v17
+ ; GCN-NEXT: ; implicit-def: $vgpr240
+ ; GCN-NEXT: v_add_u32_e32 v208, 0x80, v1
; GCN-NEXT: ; implicit-def: $vgpr156_vgpr157_vgpr158_vgpr159
; GCN-NEXT: ; implicit-def: $vgpr144_vgpr145_vgpr146_vgpr147
- ; GCN-NEXT: ; implicit-def: $vgpr19
- ; GCN-NEXT: ; implicit-def: $vgpr26
- ; GCN-NEXT: ; implicit-def: $vgpr27
- ; GCN-NEXT: v_add_u32_e32 v227, 0xc0, v17
- ; GCN-NEXT: v_add_u32_e32 v231, v19, v26
- ; GCN-NEXT: v_add_u32_e32 v232, v19, v27
- ; GCN-NEXT: ; implicit-def: $sgpr0_sgpr1_sgpr2_sgpr3
- ; GCN-NEXT: ; implicit-def: $vgpr28
- ; GCN-NEXT: ; implicit-def: $vgpr29
- ; GCN-NEXT: v_add_u32_e32 v233, v19, v28
- ; GCN-NEXT: v_add_u32_e32 v234, v19, v29
+ ; GCN-NEXT: ; implicit-def: $vgpr3
+ ; GCN-NEXT: ; implicit-def: $vgpr10
+ ; GCN-NEXT: ; implicit-def: $vgpr11
+ ; GCN-NEXT: ; implicit-def: $vgpr12
+ ; GCN-NEXT: ; implicit-def: $vgpr13
+ ; GCN-NEXT: v_add_u32_e32 v225, 0xc0, v1
+ ; GCN-NEXT: v_add_u32_e32 v226, v3, v10
+ ; GCN-NEXT: v_add_u32_e32 v227, v3, v11
+ ; GCN-NEXT: v_add_u32_e32 v228, v3, v12
+ ; GCN-NEXT: v_add_u32_e32 v232, v3, v13
; GCN-NEXT: ; implicit-def: $vgpr140_vgpr141_vgpr142_vgpr143
- ; GCN-NEXT: ; implicit-def: $sgpr5
- ; GCN-NEXT: ; implicit-def: $sgpr7
+ ; GCN-NEXT: ; implicit-def: $sgpr0_sgpr1_sgpr2_sgpr3
; GCN-NEXT: ; implicit-def: $vgpr148_vgpr149_vgpr150_vgpr151
; GCN-NEXT: ; implicit-def: $vgpr136_vgpr137_vgpr138_vgpr139
+ ; GCN-NEXT: ; implicit-def: $sgpr5
+ ; GCN-NEXT: ; implicit-def: $sgpr7
; GCN-NEXT: ; implicit-def: $vgpr132_vgpr133_vgpr134_vgpr135
- ; GCN-NEXT: ; implicit-def: $vgpr20
- ; GCN-NEXT: v_add_u32_e32 v18, s17, v20
- ; GCN-NEXT: v_and_b32_e32 v18, 0x1fffffff, v18
- ; GCN-NEXT: ; implicit-def: $sgpr16
- ; GCN-NEXT: v_mul_lo_u32 v18, v18, s16
- ; GCN-NEXT: ; implicit-def: $vgpr21
- ; GCN-NEXT: v_add_lshl_u32 v199, v21, v18, 1
- ; GCN-NEXT: ; implicit-def: $vgpr22
- ; GCN-NEXT: v_lshl_add_u32 v200, v22, 1, v199
- ; GCN-NEXT: ; implicit-def: $vgpr23
- ; GCN-NEXT: v_lshl_add_u32 v201, v23, 1, v200
- ; GCN-NEXT: ; implicit-def: $vgpr24
- ; GCN-NEXT: v_lshl_add_u32 v202, v24, 1, v201
- ; GCN-NEXT: ; implicit-def: $vgpr16
- ; GCN-NEXT: ; implicit-def: $vgpr18
- ; GCN-NEXT: ; implicit-def: $vgpr20
- ; GCN-NEXT: ; implicit-def: $vgpr24
- ; GCN-NEXT: v_add_u32_e32 v247, v19, v24
- ; GCN-NEXT: v_add_u32_e32 v248, v19, v16
- ; GCN-NEXT: v_add_u32_e32 v249, v19, v18
- ; GCN-NEXT: v_add_u32_e32 v250, v19, v20
; GCN-NEXT: ; implicit-def: $vgpr128_vgpr129_vgpr130_vgpr131
; GCN-NEXT: ; implicit-def: $sgpr14
- ; GCN-NEXT: ; implicit-def: $vgpr196
+ ; GCN-NEXT: ; implicit-def: $vgpr6
+ ; GCN-NEXT: v_lshl_add_u32 v173, v6, 1, v172
+ ; GCN-NEXT: ; implicit-def: $vgpr7
+ ; GCN-NEXT: v_lshl_add_u32 v174, v7, 1, v173
+ ; GCN-NEXT: ; implicit-def: $vgpr8
+ ; GCN-NEXT: v_lshl_add_u32 v175, v8, 1, v174
+ ; GCN-NEXT: ; implicit-def: $vgpr0
+ ; GCN-NEXT: ; implicit-def: $vgpr2
+ ; GCN-NEXT: ; implicit-def: $vgpr4
+ ; GCN-NEXT: ; implicit-def: $vgpr170
+ ; GCN-NEXT: v_add_u32_e32 v242, v3, v170
+ ; GCN-NEXT: v_add_u32_e32 v243, v3, v0
+ ; GCN-NEXT: v_add_u32_e32 v244, v3, v2
+ ; GCN-NEXT: v_add_u32_e32 v245, v3, v4
+ ; GCN-NEXT: ; implicit-def: $vgpr160
; GCN-NEXT: ; implicit-def: $sgpr12_sgpr13
- ; GCN-NEXT: ; implicit-def: $vgpr211
- ; GCN-NEXT: v_max_f32_e32 v212, v211, v211
- ; GCN-NEXT: ; implicit-def: $vgpr198
- ; GCN-NEXT: ; implicit-def: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15
- ; GCN-NEXT: ; implicit-def: $vgpr32
- ; GCN-NEXT: ; implicit-def: $vgpr33
- ; GCN-NEXT: ; implicit-def: $vgpr34
- ; GCN-NEXT: v_add_u32_e32 v210, v19, v34
- ; GCN-NEXT: v_add_u32_e32 v206, v19, v33
- ; GCN-NEXT: v_add_u32_e32 v205, v19, v32
- ; GCN-NEXT: ; implicit-def: $vgpr32_vgpr33_vgpr34_vgpr35_vgpr36_vgpr37_vgpr38_vgpr39_vgpr40_vgpr41_vgpr42_vgpr43_vgpr44_vgpr45_vgpr46_vgpr47
- ; GCN-NEXT: ; implicit-def: $vgpr21
- ; GCN-NEXT: ; implicit-def: $vgpr22
- ; GCN-NEXT: ; implicit-def: $vgpr23
- ; GCN-NEXT: ; implicit-def: $vgpr30
- ; GCN-NEXT: ; implicit-def: $vgpr31
- ; GCN-NEXT: v_add_u32_e32 v207, v19, v21
- ; GCN-NEXT: v_add_u32_e32 v208, v19, v22
- ; GCN-NEXT: v_add_u32_e32 v209, v19, v23
- ; GCN-NEXT: v_add_u32_e32 v203, v19, v30
- ; GCN-NEXT: v_add_u32_e32 v204, v19, v31
- ; GCN-NEXT: ; kill: killed $vgpr17
+ ; GCN-NEXT: ; implicit-def: $vgpr241
+ ; GCN-NEXT: v_max_f32_e32 v246, v241, v241
+ ; GCN-NEXT: ; implicit-def: $vgpr171
; GCN-NEXT: ; implicit-def: $vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31
+ ; GCN-NEXT: ; implicit-def: $vgpr32_vgpr33_vgpr34_vgpr35_vgpr36_vgpr37_vgpr38_vgpr39_vgpr40_vgpr41_vgpr42_vgpr43_vgpr44_vgpr45_vgpr46_vgpr47
+ ; GCN-NEXT: ; implicit-def: $vgpr161
+ ; GCN-NEXT: ; implicit-def: $vgpr162
+ ; GCN-NEXT: ; implicit-def: $vgpr163
+ ; GCN-NEXT: ; implicit-def: $vgpr165
+ ; GCN-NEXT: v_add_u32_e32 v247, v3, v165
+ ; GCN-NEXT: v_add_u32_e32 v248, v3, v161
+ ; GCN-NEXT: v_add_u32_e32 v249, v3, v162
+ ; GCN-NEXT: v_add_u32_e32 v250, v3, v163
+ ; GCN-NEXT: ; implicit-def: $vgpr164
+ ; GCN-NEXT: ; implicit-def: $vgpr166
+ ; GCN-NEXT: ; implicit-def: $vgpr167
+ ; GCN-NEXT: ; implicit-def: $vgpr168
+ ; GCN-NEXT: v_add_u32_e32 v176, v3, v168
+ ; GCN-NEXT: v_add_u32_e32 v177, v3, v164
+ ; GCN-NEXT: v_add_u32_e32 v178, v3, v166
+ ; GCN-NEXT: v_add_u32_e32 v179, v3, v167
+ ; GCN-NEXT: ; kill: killed $vgpr1
+ ; GCN-NEXT: ; implicit-def: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15
; GCN-NEXT: ; implicit-def: $vgpr48_vgpr49_vgpr50_vgpr51_vgpr52_vgpr53_vgpr54_vgpr55_vgpr56_vgpr57_vgpr58_vgpr59_vgpr60_vgpr61_vgpr62_vgpr63
- ; GCN-NEXT: ; implicit-def: $vgpr197
+ ; GCN-NEXT: ; implicit-def: $vgpr169
; GCN-NEXT: ; iglp_opt mask(0x00000002)
; GCN-NEXT: buffer_wbl2 sc0 sc1
- ; GCN-NEXT: ds_write_b128 v230, v[64:67]
+ ; GCN-NEXT: ds_write_b128 v231, v[64:67]
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt lgkmcnt(0)
- ; GCN-NEXT: ds_write_b128 v230, v[68:71] offset:1024
+ ; GCN-NEXT: ds_write_b128 v231, v[68:71] offset:1024
; GCN-NEXT: s_waitcnt lgkmcnt(0)
- ; GCN-NEXT: buffer_load_dwordx4 v[160:163], v226, s[8:11], 0 offen offset:64 sc0 sc1
+ ; GCN-NEXT: buffer_load_dwordx4 v[180:183], v224, s[8:11], 0 offen offset:64 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: buffer_load_dwordx4 v[164:167], v72, s[8:11], 0 offen sc0 sc1
+ ; GCN-NEXT: buffer_load_dwordx4 v[184:187], v72, s[8:11], 0 offen sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: ;;#ASMSTART
; GCN-NEXT: s_waitcnt vmcnt(8)
; GCN-NEXT: ;;#ASMEND
- ; GCN-NEXT: ds_read_b128 v[64:67], v213
+ ; GCN-NEXT: ds_read_b128 v[64:67], v230
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[64:65], v[152:153], 0
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[66:67], v[154:155], v[112:127]
- ; GCN-NEXT: ds_read_b128 v[64:67], v213 offset:512
+ ; GCN-NEXT: ds_read_b128 v[64:67], v230 offset:512
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[64:65], v[152:153], 0
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[96:111], v[66:67], v[154:155], v[96:111]
- ; GCN-NEXT: ds_read_b128 v[64:67], v213 offset:1024
+ ; GCN-NEXT: ds_read_b128 v[64:67], v230 offset:1024
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: ds_read_b128 v[168:171], v213 offset:1536
+ ; GCN-NEXT: ds_read_b128 v[188:191], v230 offset:1536
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: ds_read_b128 v[172:175], v246
+ ; GCN-NEXT: ds_read_b128 v[192:195], v240
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: ds_read_b128 v[176:179], v246 offset:512
+ ; GCN-NEXT: ds_read_b128 v[196:199], v240 offset:512
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: ds_read_b128 v[180:183], v246 offset:1024
+ ; GCN-NEXT: ds_read_b128 v[200:203], v240 offset:1024
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: ds_read_b128 v[184:187], v246 offset:1536
+ ; GCN-NEXT: ds_read_b128 v[204:207], v240 offset:1536
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: ;;#ASMSTART
@@ -150,293 +150,293 @@
; GCN-NEXT: ;;#ASMEND
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[64:65], v[152:153], 0
; GCN-NEXT: buffer_wbl2 sc0 sc1
- ; GCN-NEXT: ds_write_b128 v230, v[160:163]
+ ; GCN-NEXT: ds_write_b128 v231, v[180:183]
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[80:95], v[66:67], v[154:155], v[80:95]
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt lgkmcnt(0)
- ; GCN-NEXT: ds_write_b128 v230, v[164:167] offset:1024
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[168:169], v[152:153], 0
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[170:171], v[154:155], v[64:79]
+ ; GCN-NEXT: ds_write_b128 v231, v[184:187] offset:1024
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[188:189], v[152:153], 0
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[64:79], v[190:191], v[154:155], v[64:79]
; GCN-NEXT: s_waitcnt lgkmcnt(0)
- ; GCN-NEXT: buffer_load_dwordx4 v[152:155], v226, s[8:11], 0 offen offset:128 sc0 sc1
+ ; GCN-NEXT: buffer_load_dwordx4 v[152:155], v224, s[8:11], 0 offen offset:128 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: buffer_load_dwordx4 v[160:163], v188, s[8:11], 0 offen sc0 sc1
+ ; GCN-NEXT: buffer_load_dwordx4 v[180:183], v208, s[8:11], 0 offen sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: ;;#ASMSTART
; GCN-NEXT: s_waitcnt vmcnt(8)
; GCN-NEXT: ;;#ASMEND
- ; GCN-NEXT: ds_read_b128 v[188:191], v213
+ ; GCN-NEXT: ds_read_b128 v[184:187], v230
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: ds_read_b128 v[192:195], v213 offset:512
+ ; GCN-NEXT: ds_read_b128 v[188:191], v230 offset:512
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: ds_read_b128 v[164:167], v213 offset:1024
+ ; GCN-NEXT: ds_read_b128 v[208:211], v230 offset:1024
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: ds_read_b128 v[214:217], v213 offset:1536
+ ; GCN-NEXT: ds_read_b128 v[212:215], v230 offset:1536
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[172:173], v[156:157], v[112:127]
- ; GCN-NEXT: ds_read_b128 v[218:221], v246
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[192:193], v[156:157], v[112:127]
+ ; GCN-NEXT: ds_read_b128 v[216:219], v240
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: ds_read_b128 v[222:225], v246 offset:512
+ ; GCN-NEXT: ds_read_b128 v[220:223], v240 offset:512
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: ds_read_b128 v[168:171], v246 offset:1024
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[112:127], v[194:195], v[158:159], v[112:127]
+ ; GCN-NEXT: ds_read_b128 v[192:195], v240 offset:1024
; GCN-NEXT: s_waitcnt...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/170242
More information about the llvm-commits
mailing list