[llvm] [AMDGPU] IGLP: Fix static variables (alternative to #137549) (PR #185436)
Robert Imschweiler via llvm-commits
llvm-commits at lists.llvm.org
Mon Mar 9 07:56:45 PDT 2026
https://github.com/ro-i created https://github.com/llvm/llvm-project/pull/185436
This is an alternative PR to https://github.com/llvm/llvm-project/pull/137549 that just fixes the correctness issue in the IGLP code for now, without caching.
Assisted-by: claude-4.6-opus-high
>From 773eab355a7f44ae988c07e57674b1d2ea7208ce Mon Sep 17 00:00:00 2001
From: Robert Imschweiler <robert.imschweiler at amd.com>
Date: Wed, 26 Nov 2025 09:14:11 -0600
Subject: [PATCH 1/3] [AMDGPU] IGLP: Fix static variables
Replace global / class-level static variables with instance members and
guarantee thread-safety.
---
llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp | 210 ++++++++++++------
.../lib/Target/AMDGPU/SIMachineFunctionInfo.h | 60 ++++-
.../CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.ll | 28 +++
3 files changed, 231 insertions(+), 67 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
index e13147e9461eb..f99514b84594e 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
@@ -833,6 +833,8 @@ class IGLPStrategy {
const SIInstrInfo *TII;
+ const MachineInstr *IGLPOptMI;
+
public:
/// Add SchedGroups to \p SyncedSchedGroups to implement this Strategy.
virtual bool applyIGLPStrategy(
@@ -846,8 +848,9 @@ class IGLPStrategy {
bool IsBottomUp = true;
- IGLPStrategy(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII)
- : DAG(DAG), TII(TII) {}
+ IGLPStrategy(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII,
+ const MachineInstr *IGLPOptMI)
+ : DAG(DAG), TII(TII), IGLPOptMI(IGLPOptMI) {}
virtual ~IGLPStrategy() = default;
};
@@ -865,8 +868,9 @@ class MFMASmallGemmOpt final : public IGLPStrategy {
return true;
}
- MFMASmallGemmOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII)
- : IGLPStrategy(DAG, TII) {
+ MFMASmallGemmOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII,
+ const MachineInstr *IGLPOptMI)
+ : IGLPStrategy(DAG, TII, IGLPOptMI) {
IsBottomUp = true;
}
};
@@ -899,31 +903,36 @@ bool MFMASmallGemmOpt::applyIGLPStrategy(
class MFMAExpInterleaveOpt final : public IGLPStrategy {
private:
// The count of TRANS SUs involved in the interleaved pipeline
- static unsigned TransPipeCount;
+ unsigned TransPipeCount = 0;
// The count of MFMA SUs involved in the interleaved pipeline
- static unsigned MFMAPipeCount;
+ unsigned MFMAPipeCount = 0;
// The count of Add SUs involved in the interleaved pipeline
- static unsigned AddPipeCount;
+ unsigned AddPipeCount = 0;
// The number of transitive MFMA successors for each TRANS SU
- static unsigned MFMAEnablement;
+ unsigned MFMAEnablement = 0;
// The number of transitive TRANS predecessors for each MFMA SU
- static unsigned ExpRequirement;
+ unsigned ExpRequirement = 0;
// The count of independent "chains" of MFMA instructions in the pipeline
- static unsigned MFMAChains;
+ unsigned MFMAChains = 0;
// The length of each independent "chain" of MFMA instructions
- static unsigned MFMAChainLength;
+ unsigned MFMAChainLength = 0;
// Whether or not the pipeline has V_CVT instructions
- static bool HasCvt;
+ bool HasCvt = false;
// Whether or not there are instructions between the TRANS instruction and
// V_CVT
- static bool HasChainBetweenCvt;
+ bool HasChainBetweenCvt = false;
// The first occuring DS_READ which feeds an MFMA chain
- static std::optional<unsigned> FirstPipeDSR;
+ std::optional<unsigned> FirstPipeDSR = std::nullopt;
// The MFMAPipe SUs with no MFMA predecessors
SmallVector<SUnit *, 4> MFMAChainSeeds;
// Compute the heuristics for the pipeline, returning whether or not the DAG
// is well formatted for the mutation
- bool analyzeDAG(const SIInstrInfo *TII);
+ bool analyzeDAG(const SIInstrInfo *TII, AMDGPU::SchedulingPhase Phase);
+ bool computeDAGAnalysis(const SIInstrInfo *TII,
+ SIMachineFunctionInfo::MFMAExpInterleaveCache &Cache);
+ void initializeFromCache(
+ const SIMachineFunctionInfo::MFMAExpInterleaveCache &Cache);
+ bool AnalysisResult;
/// Whether or not the instruction is a transitive predecessor of an MFMA
/// instruction
@@ -1328,29 +1337,22 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
bool shouldApplyStrategy(ScheduleDAGInstrs *DAG,
AMDGPU::SchedulingPhase Phase) override;
- MFMAExpInterleaveOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII)
- : IGLPStrategy(DAG, TII) {
+ MFMAExpInterleaveOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII,
+ const MachineInstr *IGLPOptMI)
+ : IGLPStrategy(DAG, TII, IGLPOptMI) {
IsBottomUp = false;
}
};
-unsigned MFMAExpInterleaveOpt::TransPipeCount = 0;
-unsigned MFMAExpInterleaveOpt::MFMAPipeCount = 0;
-unsigned MFMAExpInterleaveOpt::AddPipeCount = 0;
-unsigned MFMAExpInterleaveOpt::MFMAEnablement = 0;
-unsigned MFMAExpInterleaveOpt::ExpRequirement = 0;
-unsigned MFMAExpInterleaveOpt::MFMAChains = 0;
-unsigned MFMAExpInterleaveOpt::MFMAChainLength = 0;
-bool MFMAExpInterleaveOpt::HasCvt = false;
-bool MFMAExpInterleaveOpt::HasChainBetweenCvt = false;
-std::optional<unsigned> MFMAExpInterleaveOpt::FirstPipeDSR = std::nullopt;
-
-bool MFMAExpInterleaveOpt::analyzeDAG(const SIInstrInfo *TII) {
+bool MFMAExpInterleaveOpt::computeDAGAnalysis(
+ const SIInstrInfo *TII,
+ SIMachineFunctionInfo::MFMAExpInterleaveCache &Cache) {
SmallVector<SUnit *, 10> ExpPipeCands;
SmallVector<SUnit *, 10> MFMAPipeCands;
SmallVector<SUnit *, 10> MFMAPipeSUs;
SmallVector<SUnit *, 10> PackSUs;
SmallVector<SUnit *, 10> CvtSUs;
+ const MachineInstr *FirstPipeDSRInstr = nullptr;
auto isBitPack = [](unsigned Opc) {
return Opc == AMDGPU::V_PACK_B32_F16_e64 || Opc == AMDGPU::V_PERM_B32_e64;
@@ -1367,12 +1369,14 @@ bool MFMAExpInterleaveOpt::analyzeDAG(const SIInstrInfo *TII) {
auto Opc = SU.getInstr()->getOpcode();
if (TII->isTRANS(Opc)) {
// Avoid counting a potential bonus V_EXP which all the MFMA depend on
+ // FIXME: This heuristic needs improvement/clarification!
+ // In general, the pipeline seems to look like this:
+ // fma_f32 -> exp_f32 -> cvt_f16_f32 -> v_pack_b32_f16 -> mfma_.._f16
+ // (with potential arithmetic between exp and cvt)
+ // see
+ // https://github.com/llvm/llvm-project/pull/80370#discussion_r1483660378
if (SU.Succs.size() >= 7)
continue;
- for (auto &Succ : SU.Succs) {
- if (Succ.getSUnit()->Succs.size() >= 7)
- continue;
- }
ExpPipeCands.push_back(&SU);
}
@@ -1457,6 +1461,7 @@ bool MFMAExpInterleaveOpt::analyzeDAG(const SIInstrInfo *TII) {
}
}
+ MFMAChainSeeds.clear();
MFMAChains = 0;
for (auto &MFMAPipeSU : MFMAPipeSUs) {
if (is_contained(MFMAChainSeeds, MFMAPipeSU))
@@ -1474,8 +1479,10 @@ bool MFMAExpInterleaveOpt::analyzeDAG(const SIInstrInfo *TII) {
for (auto Pred : MFMAChainSeeds[0]->Preds) {
if (TII->isDS(Pred.getSUnit()->getInstr()->getOpcode()) &&
- Pred.getSUnit()->getInstr()->mayLoad())
+ Pred.getSUnit()->getInstr()->mayLoad()) {
FirstPipeDSR = Pred.getSUnit()->NodeNum;
+ FirstPipeDSRInstr = Pred.getSUnit()->getInstr();
+ }
}
MFMAChainLength = MFMAPipeCount / MFMAChains;
@@ -1519,20 +1526,76 @@ bool MFMAExpInterleaveOpt::analyzeDAG(const SIInstrInfo *TII) {
});
ExpRequirement *= PackPredCount;
+
+ // Cache the results for later scheduling phases.
+ Cache.TransPipeCount = TransPipeCount;
+ Cache.MFMAPipeCount = MFMAPipeCount;
+ Cache.AddPipeCount = AddPipeCount;
+ Cache.MFMAEnablement = MFMAEnablement;
+ Cache.ExpRequirement = ExpRequirement;
+ Cache.MFMAChains = MFMAChains;
+ Cache.MFMAChainLength = MFMAChainLength;
+ Cache.HasCvt = HasCvt;
+ Cache.HasChainBetweenCvt = HasChainBetweenCvt;
+ Cache.FirstPipeDSRInstr = FirstPipeDSRInstr;
+ Cache.MFMAChainSeedInstrs.clear();
+ Cache.MFMAChainSeedInstrs.reserve(MFMAChainSeeds.size());
+ for (SUnit *Seed : MFMAChainSeeds)
+ Cache.MFMAChainSeedInstrs.push_back(Seed->getInstr());
+ Cache.AnalysisResult = true;
+
return true;
}
+void MFMAExpInterleaveOpt::initializeFromCache(
+ const SIMachineFunctionInfo::MFMAExpInterleaveCache &Cache) {
+ TransPipeCount = Cache.TransPipeCount;
+ MFMAPipeCount = Cache.MFMAPipeCount;
+ AddPipeCount = Cache.AddPipeCount;
+ MFMAEnablement = Cache.MFMAEnablement;
+ ExpRequirement = Cache.ExpRequirement;
+ MFMAChains = Cache.MFMAChains;
+ MFMAChainLength = Cache.MFMAChainLength;
+ HasCvt = Cache.HasCvt;
+ HasChainBetweenCvt = Cache.HasChainBetweenCvt;
+ SUnit *SU =
+ DAG->getSUnit(const_cast<MachineInstr *>(Cache.FirstPipeDSRInstr));
+ assert(SU && "FirstPipeDSRInstr instruction not found in DAG");
+ FirstPipeDSR = SU->NodeNum;
+ MFMAChainSeeds.clear();
+ for (const MachineInstr *MI : Cache.MFMAChainSeedInstrs) {
+ SUnit *SeedSU = DAG->getSUnit(const_cast<MachineInstr *>(MI));
+ assert(SeedSU && "MFMAChainSeed instruction not found in DAG");
+ MFMAChainSeeds.push_back(SeedSU);
+ }
+ AnalysisResult = Cache.AnalysisResult;
+}
+
+bool MFMAExpInterleaveOpt::analyzeDAG(const SIInstrInfo *TII,
+ AMDGPU::SchedulingPhase Phase) {
+ SIMachineFunctionInfo &MFI = *DAG->MF.getInfo<SIMachineFunctionInfo>();
+
+ if (const SIMachineFunctionInfo::MFMAExpInterleaveCache *Cache =
+ MFI.getMFMAExpInterleaveCache(IGLPOptMI)) {
+ initializeFromCache(*Cache);
+ return AnalysisResult;
+ }
+
+ assert(Phase != AMDGPU::SchedulingPhase::PostRA &&
+ "PostRA phase not expected to require analyzing DAG");
+ SIMachineFunctionInfo::MFMAExpInterleaveCache Cache;
+ AnalysisResult = computeDAGAnalysis(TII, Cache);
+ if (AnalysisResult)
+ MFI.setMFMAExpInterleaveCache(IGLPOptMI, Cache);
+ return AnalysisResult;
+}
+
bool MFMAExpInterleaveOpt::shouldApplyStrategy(ScheduleDAGInstrs *DAG,
AMDGPU::SchedulingPhase Phase) {
const GCNSubtarget &ST = DAG->MF.getSubtarget<GCNSubtarget>();
const SIInstrInfo *TII = ST.getInstrInfo();
- if (Phase != AMDGPU::SchedulingPhase::PostRA)
- MFMAChainSeeds.clear();
- if (Phase != AMDGPU::SchedulingPhase::PostRA && !analyzeDAG(TII))
- return false;
-
- return true;
+ return analyzeDAG(TII, Phase);
}
bool MFMAExpInterleaveOpt::applyIGLPStrategy(
@@ -1540,6 +1603,8 @@ bool MFMAExpInterleaveOpt::applyIGLPStrategy(
DenseMap<int, SmallVector<SchedGroup, 4>> &SyncedSchedGroups,
AMDGPU::SchedulingPhase Phase) {
+ assert(AnalysisResult && "no or failed DAG analysis");
+
bool IsSmallKernelType =
MFMAEnablement == 2 && ExpRequirement == 4 && TransPipeCount == 32;
bool IsLargeKernelType =
@@ -1559,18 +1624,18 @@ bool MFMAExpInterleaveOpt::applyIGLPStrategy(
unsigned CurrMFMAForTransPosition = 0;
auto incrementTransPosition = [&MFMAChain, &PositionInChain,
- &CurrMFMAForTransPosition]() {
+ &CurrMFMAForTransPosition, this]() {
CurrMFMAForTransPosition += MFMAEnablement;
PositionInChain = (CurrMFMAForTransPosition / MFMAChains);
MFMAChain = CurrMFMAForTransPosition % MFMAChains;
};
- auto getNextTransPositionInChain = [&CurrMFMAForTransPosition]() {
+ auto getNextTransPositionInChain = [&CurrMFMAForTransPosition, this]() {
auto TempMFMAForTrans = CurrMFMAForTransPosition + MFMAEnablement;
return (TempMFMAForTrans / MFMAChains);
};
- auto getNextTransMFMAChain = [&CurrMFMAForTransPosition]() {
+ auto getNextTransMFMAChain = [&CurrMFMAForTransPosition, this]() {
auto TempMFMAForTrans = CurrMFMAForTransPosition + MFMAEnablement;
return TempMFMAForTrans % MFMAChains;
};
@@ -1580,7 +1645,7 @@ bool MFMAExpInterleaveOpt::applyIGLPStrategy(
unsigned PositionInChainForMFMA = 0;
auto incrementMFMAPosition = [&CurrMFMAPosition, &MFMAChainForMFMA,
- &PositionInChainForMFMA]() {
+ &PositionInChainForMFMA, this]() {
++CurrMFMAPosition;
MFMAChainForMFMA = CurrMFMAPosition % MFMAChains;
PositionInChainForMFMA = CurrMFMAPosition / MFMAChains;
@@ -1838,8 +1903,9 @@ class MFMAExpSimpleInterleaveOpt final : public IGLPStrategy {
return true;
}
- MFMAExpSimpleInterleaveOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII)
- : IGLPStrategy(DAG, TII) {
+ MFMAExpSimpleInterleaveOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII,
+ const MachineInstr *IGLPOptMI)
+ : IGLPStrategy(DAG, TII, IGLPOptMI) {
IsBottomUp = true;
}
};
@@ -2065,25 +2131,36 @@ class MFMASmallGemmSingleWaveOpt final : public IGLPStrategy {
return true;
}
- MFMASmallGemmSingleWaveOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII)
- : IGLPStrategy(DAG, TII) {
+ MFMASmallGemmSingleWaveOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII,
+ const MachineInstr *IGLPOptMI)
+ : IGLPStrategy(DAG, TII, IGLPOptMI) {
IsBottomUp = false;
}
};
-static unsigned DSWCount = 0;
-static unsigned DSWWithPermCount = 0;
-static unsigned DSWWithSharedVMEMCount = 0;
-
bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
DenseMap<int, SUnitsToCandidateSGsMap> &SyncedInstrs,
DenseMap<int, SmallVector<SchedGroup, 4>> &SyncedSchedGroups,
AMDGPU::SchedulingPhase Phase) {
unsigned MFMACount = 0;
unsigned DSRCount = 0;
+ unsigned DSWCount = 0;
+ unsigned DSWWithPermCount = 0;
+ unsigned DSWWithSharedVMEMCount = 0;
bool IsInitial = Phase == AMDGPU::SchedulingPhase::Initial;
+ if (!IsInitial) {
+ const SIMachineFunctionInfo &MFI =
+ *DAG->MF.getInfo<SIMachineFunctionInfo>();
+ const SIMachineFunctionInfo::MFMASmallGemmSingleWaveCache *Cache =
+ MFI.getMFMASmallGemmSingleWaveCache(IGLPOptMI);
+ assert(Cache && "no cache found");
+ DSWCount = Cache->DSWCount;
+ DSWWithPermCount = Cache->DSWWithPermCount;
+ DSWWithSharedVMEMCount = Cache->DSWWithSharedVMEMCount;
+ }
+
assert((!IsInitial || (DSWCount == 0 && DSWWithPermCount == 0 &&
DSWWithSharedVMEMCount == 0)) &&
"DSWCounters should be zero in pre-RA scheduling!");
@@ -2110,8 +2187,6 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
if (IsInitial) {
DSWWithPermCount = DSWithPerms.size();
- auto *I = DSWithPerms.begin();
- auto *E = DSWithPerms.end();
// Get the count of DS_WRITES with V_PERM predecessors which
// have loop carried dependencies (WAR) on the same VMEM_READs.
@@ -2121,10 +2196,10 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
// for every V_PERM pred of this DS_W.
DenseMap<MachineInstr *, SUnit *> VMEMLookup;
SmallVector<SUnit *, 6> Counted;
- for (; I != E; I++) {
+ for (SUnit *DSWrite : DSWithPerms) {
SUnit *Cand = nullptr;
bool MissedAny = false;
- for (auto &Pred : (*I)->Preds) {
+ for (auto &Pred : DSWrite->Preds) {
if (Pred.getSUnit()->getInstr()->getOpcode() != AMDGPU::V_PERM_B32_e64)
continue;
@@ -2138,11 +2213,11 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
if (MissedAny || !VMEMLookup.size()) {
MissedAny = true;
- VMEMLookup[MI] = *I;
+ VMEMLookup[MI] = DSWrite;
continue;
}
- auto [It, Inserted] = VMEMLookup.try_emplace(MI, *I);
+ auto [It, Inserted] = VMEMLookup.try_emplace(MI, DSWrite);
if (Inserted) {
MissedAny = true;
continue;
@@ -2158,9 +2233,14 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
if (!MissedAny && Cand) {
DSWWithSharedVMEMCount += 2;
Counted.push_back(Cand);
- Counted.push_back(*I);
+ Counted.push_back(DSWrite);
}
}
+
+ SIMachineFunctionInfo::MFMASmallGemmSingleWaveCache Cache = {
+ DSWCount, DSWWithPermCount, DSWWithSharedVMEMCount};
+ SIMachineFunctionInfo &MFI = *DAG->MF.getInfo<SIMachineFunctionInfo>();
+ MFI.setMFMASmallGemmSingleWaveCache(IGLPOptMI, Cache);
}
assert(DSWWithSharedVMEMCount <= DSWWithPermCount);
@@ -2322,16 +2402,16 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
static std::unique_ptr<IGLPStrategy>
createIGLPStrategy(IGLPStrategyID ID, ScheduleDAGInstrs *DAG,
- const SIInstrInfo *TII) {
+ const SIInstrInfo *TII, const MachineInstr *IGLPOptMI) {
switch (ID) {
case MFMASmallGemmOptID:
- return std::make_unique<MFMASmallGemmOpt>(DAG, TII);
+ return std::make_unique<MFMASmallGemmOpt>(DAG, TII, IGLPOptMI);
case MFMASmallGemmSingleWaveOptID:
- return std::make_unique<MFMASmallGemmSingleWaveOpt>(DAG, TII);
+ return std::make_unique<MFMASmallGemmSingleWaveOpt>(DAG, TII, IGLPOptMI);
case MFMAExpInterleaveID:
- return std::make_unique<MFMAExpInterleaveOpt>(DAG, TII);
+ return std::make_unique<MFMAExpInterleaveOpt>(DAG, TII, IGLPOptMI);
case MFMAExpSimpleInterleaveID:
- return std::make_unique<MFMAExpSimpleInterleaveOpt>(DAG, TII);
+ return std::make_unique<MFMAExpSimpleInterleaveOpt>(DAG, TII, IGLPOptMI);
}
llvm_unreachable("Unknown IGLPStrategyID");
@@ -2709,7 +2789,7 @@ void IGroupLPDAGMutation::initSchedGroupBarrierPipelineStage(
bool IGroupLPDAGMutation::initIGLPOpt(SUnit &SU) {
IGLPStrategyID StrategyID =
(IGLPStrategyID)SU.getInstr()->getOperand(0).getImm();
- auto S = createIGLPStrategy(StrategyID, DAG, TII);
+ auto S = createIGLPStrategy(StrategyID, DAG, TII, SU.getInstr());
if (!S->shouldApplyStrategy(DAG, Phase))
return false;
diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
index 617862db8f506..bc1655af1fe73 100644
--- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
@@ -20,6 +20,7 @@
#include "MCTargetDesc/AMDGPUMCTargetDesc.h"
#include "SIInstrInfo.h"
#include "SIModeRegisterDefaults.h"
+#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/SetVector.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/CodeGen/MIRYamlMapping.h"
@@ -29,6 +30,8 @@
namespace llvm {
+class MachineInstr;
+
class MachineFrameInfo;
class MachineFunction;
class SIMachineFunctionInfo;
@@ -550,6 +553,27 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
bool IsDead = false;
};
+ struct MFMASmallGemmSingleWaveCache {
+ unsigned DSWCount = 0;
+ unsigned DSWWithPermCount = 0;
+ unsigned DSWWithSharedVMEMCount = 0;
+ };
+
+ struct MFMAExpInterleaveCache {
+ SmallVector<const MachineInstr *, 4> MFMAChainSeedInstrs;
+ const MachineInstr *FirstPipeDSRInstr = nullptr;
+ unsigned TransPipeCount = 0;
+ unsigned MFMAPipeCount = 0;
+ unsigned AddPipeCount = 0;
+ unsigned MFMAEnablement = 0;
+ unsigned ExpRequirement = 0;
+ unsigned MFMAChains = 0;
+ unsigned MFMAChainLength = 0;
+ bool HasCvt = false;
+ bool HasChainBetweenCvt = false;
+ bool AnalysisResult = false;
+ };
+
private:
// To track virtual VGPR + lane index for each subregister of the SGPR spilled
// to frameindex key during SILowerSGPRSpills pass.
@@ -614,7 +638,11 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
// load/store is enabled.
IndexedMap<uint32_t, VGPRBlock2IndexFunctor> MaskForVGPRBlockOps;
-private:
+ DenseMap<const MachineInstr *, MFMASmallGemmSingleWaveCache>
+ MFMASmallGemmSingleWaveCaches;
+ DenseMap<const MachineInstr *, MFMAExpInterleaveCache>
+ MFMAExpInterleaveCaches;
+
Register VGPRForAGPRCopy;
bool allocateVirtualVGPRForSGPRSpills(MachineFunction &MF, int FI,
@@ -632,6 +660,35 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
VGPRForAGPRCopy = NewVGPRForAGPRCopy;
}
+ const MFMASmallGemmSingleWaveCache *
+ getMFMASmallGemmSingleWaveCache(const MachineInstr *MI) const {
+ if (!MI)
+ return nullptr;
+ auto It = MFMASmallGemmSingleWaveCaches.find(MI);
+ return It == MFMASmallGemmSingleWaveCaches.end() ? nullptr : &It->second;
+ }
+
+ void
+ setMFMASmallGemmSingleWaveCache(const MachineInstr *MI,
+ const MFMASmallGemmSingleWaveCache &Cache) {
+ assert(MI && "MachineInstr pointer must not be null");
+ MFMASmallGemmSingleWaveCaches[MI] = Cache;
+ }
+
+ const MFMAExpInterleaveCache *
+ getMFMAExpInterleaveCache(const MachineInstr *MI) const {
+ if (!MI)
+ return nullptr;
+ auto It = MFMAExpInterleaveCaches.find(MI);
+ return It == MFMAExpInterleaveCaches.end() ? nullptr : &It->second;
+ }
+
+ void setMFMAExpInterleaveCache(const MachineInstr *MI,
+ const MFMAExpInterleaveCache &Cache) {
+ assert(MI && "MachineInstr pointer must not be null");
+ MFMAExpInterleaveCaches[MI] = Cache;
+ }
+
bool isCalleeSavedReg(const MCPhysReg *CSRegs, MCPhysReg Reg) const;
void setMaskForVGPRBlockOps(Register RegisterBlock, uint32_t Mask) {
@@ -647,7 +704,6 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
return MaskForVGPRBlockOps.inBounds(RegisterBlock);
}
-public:
SIMachineFunctionInfo(const SIMachineFunctionInfo &MFI) = default;
SIMachineFunctionInfo(const Function &F, const GCNSubtarget *STI);
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.ll
index a10c99070d8e1..64265f6f9a505 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.ll
@@ -285,6 +285,34 @@ entry:
ret void
}
+; If we run this function after test_iglp_opt_rev_mfma_gemm, we get:
+; > Assertion `(!IsInitial || (DSWCount == 0 && DSWWithPermCount == 0 &&
+; > DSWWithSharedVMEMCount == 0)) && "DSWCounters should be zero in pre-RA
+; > scheduling!"' failed.
+; This is because, previously, the counters were global static variables which
+; weren't reset.
+define amdgpu_kernel void @test_after_test_iglp_opt_rev_mfma_gemm(ptr %src, ptr addrspace(3) %dst) {
+; GCN-LABEL: test_after_test_iglp_opt_rev_mfma_gemm:
+; GCN: ; %bb.0: ; %entry
+; GCN-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x24
+; GCN-NEXT: ; iglp_opt mask(0x00000001)
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: v_pk_mov_b32 v[0:1], s[0:1], s[0:1] op_sel:[0,1]
+; GCN-NEXT: flat_load_ubyte v0, v[0:1]
+; GCN-NEXT: s_load_dword s0, s[4:5], 0x2c
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: v_mov_b32_e32 v1, s0
+; GCN-NEXT: s_waitcnt vmcnt(0)
+; GCN-NEXT: v_and_b32_e32 v0, 1, v0
+; GCN-NEXT: ds_write_b8 v1, v0
+; GCN-NEXT: s_endpgm
+entry:
+ %a = load i1, ptr %src, align 1
+ call void @llvm.amdgcn.iglp.opt(i32 1)
+ store i1 %a, ptr addrspace(3) %dst, align 1
+ ret void
+}
+
define amdgpu_kernel void @test_iglp_opt_asm_sideeffect(ptr addrspace(3) noalias %in, ptr addrspace(3) noalias %out) #0 {
; GCN-LABEL: test_iglp_opt_asm_sideeffect:
; GCN: ; %bb.0: ; %entry
>From 97e8d0da7cb02b35d80614a8b5ae172d30754da2 Mon Sep 17 00:00:00 2001
From: Robert Imschweiler <robert.imschweiler at amd.com>
Date: Mon, 23 Feb 2026 10:08:27 -0600
Subject: [PATCH 2/3] move pointer cache to GCNScheduleDAGMILive
Assisted-by: claude-4.6-opus
---
llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp | 56 +++++++-----
llvm/lib/Target/AMDGPU/GCNSchedStrategy.h | 30 +++++++
.../Target/AMDGPU/SIMachineFunctionInfo.cpp | 77 +++++++++++++++++
.../lib/Target/AMDGPU/SIMachineFunctionInfo.h | 83 +++++++++++++++---
.../AMDGPU/iglp-cache-serialization.mir | 85 +++++++++++++++++++
5 files changed, 302 insertions(+), 29 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/iglp-cache-serialization.mir
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
index f99514b84594e..29699962d324c 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
@@ -16,6 +16,7 @@
//===----------------------------------------------------------------------===//
#include "AMDGPUIGroupLP.h"
+#include "GCNSchedStrategy.h"
#include "MCTargetDesc/AMDGPUMCTargetDesc.h"
#include "SIInstrInfo.h"
#include "SIMachineFunctionInfo.h"
@@ -925,13 +926,13 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
std::optional<unsigned> FirstPipeDSR = std::nullopt;
// The MFMAPipe SUs with no MFMA predecessors
SmallVector<SUnit *, 4> MFMAChainSeeds;
- // Compute the heuristics for the pipeline, returning whether or not the DAG
- // is well formatted for the mutation
bool analyzeDAG(const SIInstrInfo *TII, AMDGPU::SchedulingPhase Phase);
bool computeDAGAnalysis(const SIInstrInfo *TII,
SIMachineFunctionInfo::MFMAExpInterleaveCache &Cache);
- void initializeFromCache(
+ void initializeScalarsFromCache(
const SIMachineFunctionInfo::MFMAExpInterleaveCache &Cache);
+ void initializePointersFromCache(
+ const GCNScheduleDAGMILive::IGLPExpInterleavePointerCache &PtrCache);
bool AnalysisResult;
/// Whether or not the instruction is a transitive predecessor of an MFMA
@@ -1527,7 +1528,6 @@ bool MFMAExpInterleaveOpt::computeDAGAnalysis(
ExpRequirement *= PackPredCount;
- // Cache the results for later scheduling phases.
Cache.TransPipeCount = TransPipeCount;
Cache.MFMAPipeCount = MFMAPipeCount;
Cache.AddPipeCount = AddPipeCount;
@@ -1537,17 +1537,20 @@ bool MFMAExpInterleaveOpt::computeDAGAnalysis(
Cache.MFMAChainLength = MFMAChainLength;
Cache.HasCvt = HasCvt;
Cache.HasChainBetweenCvt = HasChainBetweenCvt;
- Cache.FirstPipeDSRInstr = FirstPipeDSRInstr;
- Cache.MFMAChainSeedInstrs.clear();
- Cache.MFMAChainSeedInstrs.reserve(MFMAChainSeeds.size());
- for (SUnit *Seed : MFMAChainSeeds)
- Cache.MFMAChainSeedInstrs.push_back(Seed->getInstr());
Cache.AnalysisResult = true;
+ GCNScheduleDAGMILive *GCNDAG = static_cast<GCNScheduleDAGMILive *>(DAG);
+ GCNScheduleDAGMILive::IGLPExpInterleavePointerCache PtrCache;
+ PtrCache.FirstPipeDSRInstr = FirstPipeDSRInstr;
+ PtrCache.MFMAChainSeedInstrs.reserve(MFMAChainSeeds.size());
+ for (SUnit *Seed : MFMAChainSeeds)
+ PtrCache.MFMAChainSeedInstrs.push_back(Seed->getInstr());
+ GCNDAG->setIGLPExpInterleavePointerCache(IGLPOptMI, PtrCache);
+
return true;
}
-void MFMAExpInterleaveOpt::initializeFromCache(
+void MFMAExpInterleaveOpt::initializeScalarsFromCache(
const SIMachineFunctionInfo::MFMAExpInterleaveCache &Cache) {
TransPipeCount = Cache.TransPipeCount;
MFMAPipeCount = Cache.MFMAPipeCount;
@@ -1558,31 +1561,46 @@ void MFMAExpInterleaveOpt::initializeFromCache(
MFMAChainLength = Cache.MFMAChainLength;
HasCvt = Cache.HasCvt;
HasChainBetweenCvt = Cache.HasChainBetweenCvt;
- SUnit *SU =
- DAG->getSUnit(const_cast<MachineInstr *>(Cache.FirstPipeDSRInstr));
- assert(SU && "FirstPipeDSRInstr instruction not found in DAG");
- FirstPipeDSR = SU->NodeNum;
+ AnalysisResult = Cache.AnalysisResult;
+}
+
+void MFMAExpInterleaveOpt::initializePointersFromCache(
+ const GCNScheduleDAGMILive::IGLPExpInterleavePointerCache &PtrCache) {
+ if (PtrCache.FirstPipeDSRInstr) {
+ SUnit *SU =
+ DAG->getSUnit(const_cast<MachineInstr *>(PtrCache.FirstPipeDSRInstr));
+ assert(SU && "FirstPipeDSRInstr instruction not found in DAG");
+ FirstPipeDSR = SU->NodeNum;
+ }
MFMAChainSeeds.clear();
- for (const MachineInstr *MI : Cache.MFMAChainSeedInstrs) {
+ for (const MachineInstr *MI : PtrCache.MFMAChainSeedInstrs) {
SUnit *SeedSU = DAG->getSUnit(const_cast<MachineInstr *>(MI));
assert(SeedSU && "MFMAChainSeed instruction not found in DAG");
MFMAChainSeeds.push_back(SeedSU);
}
- AnalysisResult = Cache.AnalysisResult;
}
bool MFMAExpInterleaveOpt::analyzeDAG(const SIInstrInfo *TII,
AMDGPU::SchedulingPhase Phase) {
SIMachineFunctionInfo &MFI = *DAG->MF.getInfo<SIMachineFunctionInfo>();
+ bool IsPostRA = Phase == AMDGPU::SchedulingPhase::PostRA;
if (const SIMachineFunctionInfo::MFMAExpInterleaveCache *Cache =
MFI.getMFMAExpInterleaveCache(IGLPOptMI)) {
- initializeFromCache(*Cache);
+ initializeScalarsFromCache(*Cache);
+
+ if (!IsPostRA) {
+ GCNScheduleDAGMILive *GCNDAG = static_cast<GCNScheduleDAGMILive *>(DAG);
+ const GCNScheduleDAGMILive::IGLPExpInterleavePointerCache *PtrCache =
+ GCNDAG->getIGLPExpInterleavePointerCache(IGLPOptMI);
+ assert(PtrCache &&
+ "Pre-RA phase expected pointer cache in GCNScheduleDAGMILive");
+ initializePointersFromCache(*PtrCache);
+ }
return AnalysisResult;
}
- assert(Phase != AMDGPU::SchedulingPhase::PostRA &&
- "PostRA phase not expected to require analyzing DAG");
+ assert(!IsPostRA && "PostRA phase not expected to require analyzing DAG");
SIMachineFunctionInfo::MFMAExpInterleaveCache Cache;
AnalysisResult = computeDAGAnalysis(TII, Cache);
if (AnalysisResult)
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
index 99fd55db33285..cb07f2da679c1 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
@@ -249,6 +249,24 @@ class GCNScheduleDAGMILive final : public ScheduleDAGMILive {
friend class ILPInitialScheduleStage;
friend class RegionPressureMap;
+public:
+ struct IGLPExpInterleavePointerCache {
+ /// MFMA instructions in the interleave pipeline that have no MFMA
+ /// predecessors, i.e. the roots of independent MFMA chains.
+ SmallVector<const MachineInstr *, 4> MFMAChainSeedInstrs;
+ /// The first DS_READ that feeds into an MFMA chain seed.
+ const MachineInstr *FirstPipeDSRInstr = nullptr;
+ };
+
+private:
+ /// Pass-local cache of pointer-based IGLP analysis data, keyed by the
+ /// IGLP_OPT MachineInstr. This is safe because the DAG object lives for
+ /// the duration of the pre-RA scheduling pass and the MachineInstrs are
+ /// stable within that lifetime. Unlike SIMachineFunctionInfo caches, this
+ /// is never serialized.
+ DenseMap<const MachineInstr *, IGLPExpInterleavePointerCache>
+ IGLPExpInterleavePtrCaches;
+
const GCNSubtarget &ST;
SIMachineFunctionInfo &MFI;
@@ -317,6 +335,18 @@ class GCNScheduleDAGMILive final : public ScheduleDAGMILive {
GCNScheduleDAGMILive(MachineSchedContext *C,
std::unique_ptr<MachineSchedStrategy> S);
+ const IGLPExpInterleavePointerCache *
+ getIGLPExpInterleavePointerCache(const MachineInstr *MI) const {
+ auto It = IGLPExpInterleavePtrCaches.find(MI);
+ return It == IGLPExpInterleavePtrCaches.end() ? nullptr : &It->second;
+ }
+
+ void
+ setIGLPExpInterleavePointerCache(const MachineInstr *MI,
+ const IGLPExpInterleavePointerCache &Cache) {
+ IGLPExpInterleavePtrCaches[MI] = Cache;
+ }
+
void schedule() override;
void finalizeSchedule() override;
diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
index 32efae69b20c8..95867d442d278 100644
--- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
@@ -772,6 +772,41 @@ yaml::SIMachineFunctionInfo::SIMachineFunctionInfo(
auto SFI = MFI.getOptionalScavengeFI();
if (SFI)
ScavengeFI = yaml::FrameIndex(*SFI, MF.getFrameInfo());
+
+ const DenseMap<const MachineInstr *,
+ llvm::SIMachineFunctionInfo::MFMASmallGemmSingleWaveCache>
+ &SmallGemmCaches = MFI.getMFMASmallGemmSingleWaveCaches();
+ for (DenseMap<const MachineInstr *,
+ llvm::SIMachineFunctionInfo::MFMASmallGemmSingleWaveCache>::
+ const_iterator It = SmallGemmCaches.begin(),
+ E = SmallGemmCaches.end();
+ It != E; ++It) {
+ const MachineInstr *MI = It->first;
+ IGLPSmallGemmSingleWaveCaches.push_back(
+ {static_cast<unsigned>(MI->getParent()->getNumber()),
+ static_cast<unsigned>(MI->getOperand(0).getImm()), It->second.DSWCount,
+ It->second.DSWWithPermCount, It->second.DSWWithSharedVMEMCount});
+ }
+
+ const DenseMap<const MachineInstr *,
+ llvm::SIMachineFunctionInfo::MFMAExpInterleaveCache>
+ &ExpInterleaveCaches = MFI.getMFMAExpInterleaveCaches();
+ for (DenseMap<
+ const MachineInstr *,
+ llvm::SIMachineFunctionInfo::MFMAExpInterleaveCache>::const_iterator
+ It = ExpInterleaveCaches.begin(),
+ E = ExpInterleaveCaches.end();
+ It != E; ++It) {
+ const MachineInstr *MI = It->first;
+ IGLPExpInterleaveCaches.push_back(
+ {static_cast<unsigned>(MI->getParent()->getNumber()),
+ static_cast<unsigned>(MI->getOperand(0).getImm()),
+ It->second.TransPipeCount, It->second.MFMAPipeCount,
+ It->second.AddPipeCount, It->second.MFMAEnablement,
+ It->second.ExpRequirement, It->second.MFMAChains,
+ It->second.MFMAChainLength, It->second.HasCvt,
+ It->second.HasChainBetweenCvt, It->second.AnalysisResult});
+ }
}
void yaml::SIMachineFunctionInfo::mappingImpl(yaml::IO &YamlIO) {
@@ -821,6 +856,48 @@ bool SIMachineFunctionInfo::initializeBaseYamlFields(
} else {
ScavengeFI = std::nullopt;
}
+
+ if (!YamlMFI.IGLPSmallGemmSingleWaveCaches.empty() ||
+ !YamlMFI.IGLPExpInterleaveCaches.empty()) {
+ for (const MachineBasicBlock &MBB : MF) {
+ for (const MachineInstr &MI : MBB) {
+ if (MI.getOpcode() != AMDGPU::IGLP_OPT)
+ continue;
+ unsigned MBBNum = MBB.getNumber();
+ unsigned StrategyId = static_cast<unsigned>(MI.getOperand(0).getImm());
+
+ for (const yaml::IGLPSmallGemmSingleWaveCacheEntry &E :
+ YamlMFI.IGLPSmallGemmSingleWaveCaches) {
+ if (E.MBBNum == MBBNum && E.StrategyId == StrategyId) {
+ MFMASmallGemmSingleWaveCache Cache = {
+ E.DSWCount, E.DSWWithPermCount, E.DSWWithSharedVMEMCount};
+ MFMASmallGemmSingleWaveCaches[&MI] = Cache;
+ break;
+ }
+ }
+
+ for (const yaml::IGLPExpInterleaveCacheEntry &E :
+ YamlMFI.IGLPExpInterleaveCaches) {
+ if (E.MBBNum == MBBNum && E.StrategyId == StrategyId) {
+ MFMAExpInterleaveCache Cache;
+ Cache.TransPipeCount = E.TransPipeCount;
+ Cache.MFMAPipeCount = E.MFMAPipeCount;
+ Cache.AddPipeCount = E.AddPipeCount;
+ Cache.MFMAEnablement = E.MFMAEnablement;
+ Cache.ExpRequirement = E.ExpRequirement;
+ Cache.MFMAChains = E.MFMAChains;
+ Cache.MFMAChainLength = E.MFMAChainLength;
+ Cache.HasCvt = E.HasCvt;
+ Cache.HasChainBetweenCvt = E.HasChainBetweenCvt;
+ Cache.AnalysisResult = E.AnalysisResult;
+ MFMAExpInterleaveCaches[&MI] = Cache;
+ break;
+ }
+ }
+ }
+ }
+ }
+
return false;
}
diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
index bc1655af1fe73..7a2ac80c002d1 100644
--- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
@@ -262,6 +262,56 @@ template <> struct MappingTraits<SIMode> {
}
};
+struct IGLPSmallGemmSingleWaveCacheEntry {
+ unsigned MBBNum = 0;
+ unsigned StrategyId = 0;
+ unsigned DSWCount = 0;
+ unsigned DSWWithPermCount = 0;
+ unsigned DSWWithSharedVMEMCount = 0;
+};
+
+template <> struct MappingTraits<IGLPSmallGemmSingleWaveCacheEntry> {
+ static void mapping(IO &YamlIO, IGLPSmallGemmSingleWaveCacheEntry &E) {
+ YamlIO.mapRequired("mbb", E.MBBNum);
+ YamlIO.mapRequired("strategyId", E.StrategyId);
+ YamlIO.mapRequired("dswCount", E.DSWCount);
+ YamlIO.mapRequired("dswWithPermCount", E.DSWWithPermCount);
+ YamlIO.mapRequired("dswWithSharedVMEMCount", E.DSWWithSharedVMEMCount);
+ }
+};
+
+struct IGLPExpInterleaveCacheEntry {
+ unsigned MBBNum = 0;
+ unsigned StrategyId = 0;
+ unsigned TransPipeCount = 0;
+ unsigned MFMAPipeCount = 0;
+ unsigned AddPipeCount = 0;
+ unsigned MFMAEnablement = 0;
+ unsigned ExpRequirement = 0;
+ unsigned MFMAChains = 0;
+ unsigned MFMAChainLength = 0;
+ bool HasCvt = false;
+ bool HasChainBetweenCvt = false;
+ bool AnalysisResult = false;
+};
+
+template <> struct MappingTraits<IGLPExpInterleaveCacheEntry> {
+ static void mapping(IO &YamlIO, IGLPExpInterleaveCacheEntry &E) {
+ YamlIO.mapRequired("mbb", E.MBBNum);
+ YamlIO.mapRequired("strategyId", E.StrategyId);
+ YamlIO.mapRequired("transPipeCount", E.TransPipeCount);
+ YamlIO.mapRequired("mfmaPipeCount", E.MFMAPipeCount);
+ YamlIO.mapRequired("addPipeCount", E.AddPipeCount);
+ YamlIO.mapRequired("mfmaEnablement", E.MFMAEnablement);
+ YamlIO.mapRequired("expRequirement", E.ExpRequirement);
+ YamlIO.mapRequired("mfmaChains", E.MFMAChains);
+ YamlIO.mapRequired("mfmaChainLength", E.MFMAChainLength);
+ YamlIO.mapRequired("hasCvt", E.HasCvt);
+ YamlIO.mapRequired("hasChainBetweenCvt", E.HasChainBetweenCvt);
+ YamlIO.mapRequired("analysisResult", E.AnalysisResult);
+ }
+};
+
struct SIMachineFunctionInfo final : public yaml::MachineFunctionInfo {
uint64_t ExplicitKernArgSize = 0;
Align MaxKernArgAlign;
@@ -306,6 +356,9 @@ struct SIMachineFunctionInfo final : public yaml::MachineFunctionInfo {
bool HasInitWholeWave = false;
bool IsWholeWaveFunction = false;
+ SmallVector<IGLPSmallGemmSingleWaveCacheEntry> IGLPSmallGemmSingleWaveCaches;
+ SmallVector<IGLPExpInterleaveCacheEntry> IGLPExpInterleaveCaches;
+
unsigned DynamicVGPRBlockSize = 0;
unsigned ScratchReservedForDynamicVGPRs = 0;
@@ -363,6 +416,9 @@ template <> struct MappingTraits<SIMachineFunctionInfo> {
YamlIO.mapOptional("longBranchReservedReg", MFI.LongBranchReservedReg,
StringValue());
YamlIO.mapOptional("hasInitWholeWave", MFI.HasInitWholeWave, false);
+ YamlIO.mapOptional("iglpSmallGemmSingleWaveCaches",
+ MFI.IGLPSmallGemmSingleWaveCaches);
+ YamlIO.mapOptional("iglpExpInterleaveCaches", MFI.IGLPExpInterleaveCaches);
YamlIO.mapOptional("dynamicVGPRBlockSize", MFI.DynamicVGPRBlockSize, false);
YamlIO.mapOptional("scratchReservedForDynamicVGPRs",
MFI.ScratchReservedForDynamicVGPRs, 0);
@@ -560,8 +616,6 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
};
struct MFMAExpInterleaveCache {
- SmallVector<const MachineInstr *, 4> MFMAChainSeedInstrs;
- const MachineInstr *FirstPipeDSRInstr = nullptr;
unsigned TransPipeCount = 0;
unsigned MFMAPipeCount = 0;
unsigned AddPipeCount = 0;
@@ -662,33 +716,39 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
const MFMASmallGemmSingleWaveCache *
getMFMASmallGemmSingleWaveCache(const MachineInstr *MI) const {
- if (!MI)
- return nullptr;
- auto It = MFMASmallGemmSingleWaveCaches.find(MI);
+ DenseMap<const MachineInstr *, MFMASmallGemmSingleWaveCache>::const_iterator
+ It = MFMASmallGemmSingleWaveCaches.find(MI);
return It == MFMASmallGemmSingleWaveCaches.end() ? nullptr : &It->second;
}
void
setMFMASmallGemmSingleWaveCache(const MachineInstr *MI,
const MFMASmallGemmSingleWaveCache &Cache) {
- assert(MI && "MachineInstr pointer must not be null");
MFMASmallGemmSingleWaveCaches[MI] = Cache;
}
const MFMAExpInterleaveCache *
getMFMAExpInterleaveCache(const MachineInstr *MI) const {
- if (!MI)
- return nullptr;
- auto It = MFMAExpInterleaveCaches.find(MI);
+ DenseMap<const MachineInstr *, MFMAExpInterleaveCache>::const_iterator It =
+ MFMAExpInterleaveCaches.find(MI);
return It == MFMAExpInterleaveCaches.end() ? nullptr : &It->second;
}
void setMFMAExpInterleaveCache(const MachineInstr *MI,
const MFMAExpInterleaveCache &Cache) {
- assert(MI && "MachineInstr pointer must not be null");
MFMAExpInterleaveCaches[MI] = Cache;
}
+ const DenseMap<const MachineInstr *, MFMASmallGemmSingleWaveCache> &
+ getMFMASmallGemmSingleWaveCaches() const {
+ return MFMASmallGemmSingleWaveCaches;
+ }
+
+ const DenseMap<const MachineInstr *, MFMAExpInterleaveCache> &
+ getMFMAExpInterleaveCaches() const {
+ return MFMAExpInterleaveCaches;
+ }
+
bool isCalleeSavedReg(const MCPhysReg *CSRegs, MCPhysReg Reg) const;
void setMaskForVGPRBlockOps(Register RegisterBlock, uint32_t Mask) {
@@ -1285,4 +1345,7 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
} // end namespace llvm
+LLVM_YAML_IS_SEQUENCE_VECTOR(llvm::yaml::IGLPSmallGemmSingleWaveCacheEntry)
+LLVM_YAML_IS_SEQUENCE_VECTOR(llvm::yaml::IGLPExpInterleaveCacheEntry)
+
#endif // LLVM_LIB_TARGET_AMDGPU_SIMACHINEFUNCTIONINFO_H
diff --git a/llvm/test/CodeGen/AMDGPU/iglp-cache-serialization.mir b/llvm/test/CodeGen/AMDGPU/iglp-cache-serialization.mir
new file mode 100644
index 0000000000000..501eb308ffc48
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/iglp-cache-serialization.mir
@@ -0,0 +1,85 @@
+# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -run-pass=none -o - %s | FileCheck -check-prefix=ROUNDTRIP %s
+# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -run-pass=postmisched -o - %s | FileCheck -check-prefix=POSTRA %s
+
+# Test that IGLP cache data serializes/deserializes through MIR and is
+# available for PostRA scheduling.
+
+# Verify round-trip of the serialized cache fields.
+# ROUNDTRIP: iglpSmallGemmSingleWaveCaches:
+# ROUNDTRIP-NEXT: - mbb: 1
+# ROUNDTRIP-NEXT: strategyId: 1
+# ROUNDTRIP-NEXT: dswCount: 6
+# ROUNDTRIP-NEXT: dswWithPermCount: 2
+# ROUNDTRIP-NEXT: dswWithSharedVMEMCount: 0
+# ROUNDTRIP: iglpExpInterleaveCaches:
+# ROUNDTRIP-NEXT: - mbb: 0
+# ROUNDTRIP-NEXT: strategyId: 2
+# ROUNDTRIP-NEXT: transPipeCount: 32
+# ROUNDTRIP-NEXT: mfmaPipeCount: 40
+# ROUNDTRIP-NEXT: addPipeCount: 33
+# ROUNDTRIP-NEXT: mfmaEnablement: 2
+# ROUNDTRIP-NEXT: expRequirement: 4
+# ROUNDTRIP-NEXT: mfmaChains: 2
+# ROUNDTRIP-NEXT: mfmaChainLength: 20
+# ROUNDTRIP-NEXT: hasCvt: true
+# ROUNDTRIP-NEXT: hasChainBetweenCvt: false
+# ROUNDTRIP-NEXT: analysisResult: true
+
+# Verify PostRA scheduling doesn't crash and passes through the IGLP_OPTs.
+# POSTRA: IGLP_OPT 2
+# POSTRA: V_MFMA
+# POSTRA: IGLP_OPT 1
+--- |
+ define amdgpu_kernel void @iglp_cache_test() #0 { ret void }
+
+ attributes #0 = { "amdgpu-flat-work-group-size"="256,256" }
+...
+---
+name: iglp_cache_test
+alignment: 1
+tracksRegLiveness: true
+noPhis: true
+isSSA: false
+noVRegs: true
+machineFunctionInfo:
+ isEntryFunction: false
+ occupancy: 5
+ iglpExpInterleaveCaches:
+ - mbb: 0
+ strategyId: 2
+ transPipeCount: 32
+ mfmaPipeCount: 40
+ addPipeCount: 33
+ mfmaEnablement: 2
+ expRequirement: 4
+ mfmaChains: 2
+ mfmaChainLength: 20
+ hasCvt: true
+ hasChainBetweenCvt: false
+ analysisResult: true
+ iglpSmallGemmSingleWaveCaches:
+ - mbb: 1
+ strategyId: 1
+ dswCount: 6
+ dswWithPermCount: 2
+ dswWithSharedVMEMCount: 0
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr36_vgpr37, $vgpr40_vgpr41, $vgpr44_vgpr45, $vgpr38_vgpr39, $vgpr42_vgpr43
+
+ IGLP_OPT 2
+ $vgpr50 = V_EXP_F32_e32 $vgpr0, implicit $mode, implicit $exec
+ $vgpr51 = V_EXP_F32_e32 $vgpr1, implicit $mode, implicit $exec
+ $vgpr52 = V_ADD_F32_e32 $vgpr2, $vgpr3, implicit $mode, implicit $exec
+ early-clobber $vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31 = contract V_MFMA_F32_32X32X8F16_vgprcd_e64 $vgpr36_vgpr37, $vgpr40_vgpr41, 0, 0, 0, 0, implicit $mode, implicit $exec
+ early-clobber $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15 = contract V_MFMA_F32_32X32X8F16_vgprcd_e64 $vgpr44_vgpr45, $vgpr40_vgpr41, 0, 0, 0, 0, implicit $mode, implicit $exec
+ S_BRANCH %bb.1
+
+ bb.1:
+ liveins: $vgpr0, $vgpr1
+
+ IGLP_OPT 1
+ $vgpr0 = V_ADD_F32_e32 $vgpr0, $vgpr1, implicit $mode, implicit $exec
+ S_ENDPGM 0
+
+...
>From dec232f2798c0b6617844506e78ccedc8e28d5ef Mon Sep 17 00:00:00 2001
From: Robert Imschweiler <robert.imschweiler at amd.com>
Date: Mon, 9 Mar 2026 06:12:50 -0500
Subject: [PATCH 3/3] fix by recomputation
---
llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp | 250 ++++--------
llvm/lib/Target/AMDGPU/GCNSchedStrategy.h | 30 --
.../Target/AMDGPU/SIMachineFunctionInfo.cpp | 77 ----
.../lib/Target/AMDGPU/SIMachineFunctionInfo.h | 123 +-----
.../AMDGPU/iglp-cache-serialization.mir | 85 ----
...llvm.amdgcn.iglp.AFLCustomIRMutator.opt.ll | 14 +-
.../AMDGPU/llvm.amdgcn.iglp.opt.exp.large.mir | 370 +++++++++---------
.../AMDGPU/llvm.amdgcn.iglp.opt.exp.small.mir | 166 ++++----
8 files changed, 346 insertions(+), 769 deletions(-)
delete mode 100644 llvm/test/CodeGen/AMDGPU/iglp-cache-serialization.mir
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
index 29699962d324c..841b820a7ac42 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
@@ -16,10 +16,9 @@
//===----------------------------------------------------------------------===//
#include "AMDGPUIGroupLP.h"
-#include "GCNSchedStrategy.h"
+#include "GCNSubtarget.h"
#include "MCTargetDesc/AMDGPUMCTargetDesc.h"
#include "SIInstrInfo.h"
-#include "SIMachineFunctionInfo.h"
#include "llvm/ADT/BitmaskEnum.h"
#include "llvm/ADT/DenseMap.h"
#include "llvm/CodeGen/MachineScheduler.h"
@@ -834,8 +833,6 @@ class IGLPStrategy {
const SIInstrInfo *TII;
- const MachineInstr *IGLPOptMI;
-
public:
/// Add SchedGroups to \p SyncedSchedGroups to implement this Strategy.
virtual bool applyIGLPStrategy(
@@ -849,9 +846,8 @@ class IGLPStrategy {
bool IsBottomUp = true;
- IGLPStrategy(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII,
- const MachineInstr *IGLPOptMI)
- : DAG(DAG), TII(TII), IGLPOptMI(IGLPOptMI) {}
+ IGLPStrategy(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII)
+ : DAG(DAG), TII(TII) {}
virtual ~IGLPStrategy() = default;
};
@@ -869,9 +865,8 @@ class MFMASmallGemmOpt final : public IGLPStrategy {
return true;
}
- MFMASmallGemmOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII,
- const MachineInstr *IGLPOptMI)
- : IGLPStrategy(DAG, TII, IGLPOptMI) {
+ MFMASmallGemmOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII)
+ : IGLPStrategy(DAG, TII) {
IsBottomUp = true;
}
};
@@ -926,13 +921,9 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
std::optional<unsigned> FirstPipeDSR = std::nullopt;
// The MFMAPipe SUs with no MFMA predecessors
SmallVector<SUnit *, 4> MFMAChainSeeds;
- bool analyzeDAG(const SIInstrInfo *TII, AMDGPU::SchedulingPhase Phase);
- bool computeDAGAnalysis(const SIInstrInfo *TII,
- SIMachineFunctionInfo::MFMAExpInterleaveCache &Cache);
- void initializeScalarsFromCache(
- const SIMachineFunctionInfo::MFMAExpInterleaveCache &Cache);
- void initializePointersFromCache(
- const GCNScheduleDAGMILive::IGLPExpInterleavePointerCache &PtrCache);
+ // Compute the heuristics for the pipeline, returning whether or not the DAG
+ // is well formatted for the mutation
+ bool analyzeDAG(const SIInstrInfo *TII);
bool AnalysisResult;
/// Whether or not the instruction is a transitive predecessor of an MFMA
@@ -1338,22 +1329,18 @@ class MFMAExpInterleaveOpt final : public IGLPStrategy {
bool shouldApplyStrategy(ScheduleDAGInstrs *DAG,
AMDGPU::SchedulingPhase Phase) override;
- MFMAExpInterleaveOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII,
- const MachineInstr *IGLPOptMI)
- : IGLPStrategy(DAG, TII, IGLPOptMI) {
+ MFMAExpInterleaveOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII)
+ : IGLPStrategy(DAG, TII) {
IsBottomUp = false;
}
};
-bool MFMAExpInterleaveOpt::computeDAGAnalysis(
- const SIInstrInfo *TII,
- SIMachineFunctionInfo::MFMAExpInterleaveCache &Cache) {
+bool MFMAExpInterleaveOpt::analyzeDAG(const SIInstrInfo *TII) {
SmallVector<SUnit *, 10> ExpPipeCands;
SmallVector<SUnit *, 10> MFMAPipeCands;
SmallVector<SUnit *, 10> MFMAPipeSUs;
SmallVector<SUnit *, 10> PackSUs;
SmallVector<SUnit *, 10> CvtSUs;
- const MachineInstr *FirstPipeDSRInstr = nullptr;
auto isBitPack = [](unsigned Opc) {
return Opc == AMDGPU::V_PACK_B32_F16_e64 || Opc == AMDGPU::V_PERM_B32_e64;
@@ -1378,6 +1365,10 @@ bool MFMAExpInterleaveOpt::computeDAGAnalysis(
// https://github.com/llvm/llvm-project/pull/80370#discussion_r1483660378
if (SU.Succs.size() >= 7)
continue;
+ for (auto &Succ : SU.Succs) {
+ if (Succ.getSUnit()->Succs.size() >= 7)
+ continue;
+ }
ExpPipeCands.push_back(&SU);
}
@@ -1482,7 +1473,6 @@ bool MFMAExpInterleaveOpt::computeDAGAnalysis(
if (TII->isDS(Pred.getSUnit()->getInstr()->getOpcode()) &&
Pred.getSUnit()->getInstr()->mayLoad()) {
FirstPipeDSR = Pred.getSUnit()->NodeNum;
- FirstPipeDSRInstr = Pred.getSUnit()->getInstr();
}
}
@@ -1527,93 +1517,16 @@ bool MFMAExpInterleaveOpt::computeDAGAnalysis(
});
ExpRequirement *= PackPredCount;
-
- Cache.TransPipeCount = TransPipeCount;
- Cache.MFMAPipeCount = MFMAPipeCount;
- Cache.AddPipeCount = AddPipeCount;
- Cache.MFMAEnablement = MFMAEnablement;
- Cache.ExpRequirement = ExpRequirement;
- Cache.MFMAChains = MFMAChains;
- Cache.MFMAChainLength = MFMAChainLength;
- Cache.HasCvt = HasCvt;
- Cache.HasChainBetweenCvt = HasChainBetweenCvt;
- Cache.AnalysisResult = true;
-
- GCNScheduleDAGMILive *GCNDAG = static_cast<GCNScheduleDAGMILive *>(DAG);
- GCNScheduleDAGMILive::IGLPExpInterleavePointerCache PtrCache;
- PtrCache.FirstPipeDSRInstr = FirstPipeDSRInstr;
- PtrCache.MFMAChainSeedInstrs.reserve(MFMAChainSeeds.size());
- for (SUnit *Seed : MFMAChainSeeds)
- PtrCache.MFMAChainSeedInstrs.push_back(Seed->getInstr());
- GCNDAG->setIGLPExpInterleavePointerCache(IGLPOptMI, PtrCache);
-
return true;
}
-void MFMAExpInterleaveOpt::initializeScalarsFromCache(
- const SIMachineFunctionInfo::MFMAExpInterleaveCache &Cache) {
- TransPipeCount = Cache.TransPipeCount;
- MFMAPipeCount = Cache.MFMAPipeCount;
- AddPipeCount = Cache.AddPipeCount;
- MFMAEnablement = Cache.MFMAEnablement;
- ExpRequirement = Cache.ExpRequirement;
- MFMAChains = Cache.MFMAChains;
- MFMAChainLength = Cache.MFMAChainLength;
- HasCvt = Cache.HasCvt;
- HasChainBetweenCvt = Cache.HasChainBetweenCvt;
- AnalysisResult = Cache.AnalysisResult;
-}
-
-void MFMAExpInterleaveOpt::initializePointersFromCache(
- const GCNScheduleDAGMILive::IGLPExpInterleavePointerCache &PtrCache) {
- if (PtrCache.FirstPipeDSRInstr) {
- SUnit *SU =
- DAG->getSUnit(const_cast<MachineInstr *>(PtrCache.FirstPipeDSRInstr));
- assert(SU && "FirstPipeDSRInstr instruction not found in DAG");
- FirstPipeDSR = SU->NodeNum;
- }
- MFMAChainSeeds.clear();
- for (const MachineInstr *MI : PtrCache.MFMAChainSeedInstrs) {
- SUnit *SeedSU = DAG->getSUnit(const_cast<MachineInstr *>(MI));
- assert(SeedSU && "MFMAChainSeed instruction not found in DAG");
- MFMAChainSeeds.push_back(SeedSU);
- }
-}
-
-bool MFMAExpInterleaveOpt::analyzeDAG(const SIInstrInfo *TII,
- AMDGPU::SchedulingPhase Phase) {
- SIMachineFunctionInfo &MFI = *DAG->MF.getInfo<SIMachineFunctionInfo>();
- bool IsPostRA = Phase == AMDGPU::SchedulingPhase::PostRA;
-
- if (const SIMachineFunctionInfo::MFMAExpInterleaveCache *Cache =
- MFI.getMFMAExpInterleaveCache(IGLPOptMI)) {
- initializeScalarsFromCache(*Cache);
-
- if (!IsPostRA) {
- GCNScheduleDAGMILive *GCNDAG = static_cast<GCNScheduleDAGMILive *>(DAG);
- const GCNScheduleDAGMILive::IGLPExpInterleavePointerCache *PtrCache =
- GCNDAG->getIGLPExpInterleavePointerCache(IGLPOptMI);
- assert(PtrCache &&
- "Pre-RA phase expected pointer cache in GCNScheduleDAGMILive");
- initializePointersFromCache(*PtrCache);
- }
- return AnalysisResult;
- }
-
- assert(!IsPostRA && "PostRA phase not expected to require analyzing DAG");
- SIMachineFunctionInfo::MFMAExpInterleaveCache Cache;
- AnalysisResult = computeDAGAnalysis(TII, Cache);
- if (AnalysisResult)
- MFI.setMFMAExpInterleaveCache(IGLPOptMI, Cache);
- return AnalysisResult;
-}
-
bool MFMAExpInterleaveOpt::shouldApplyStrategy(ScheduleDAGInstrs *DAG,
AMDGPU::SchedulingPhase Phase) {
const GCNSubtarget &ST = DAG->MF.getSubtarget<GCNSubtarget>();
const SIInstrInfo *TII = ST.getInstrInfo();
- return analyzeDAG(TII, Phase);
+ AnalysisResult = analyzeDAG(TII);
+ return AnalysisResult;
}
bool MFMAExpInterleaveOpt::applyIGLPStrategy(
@@ -1921,9 +1834,8 @@ class MFMAExpSimpleInterleaveOpt final : public IGLPStrategy {
return true;
}
- MFMAExpSimpleInterleaveOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII,
- const MachineInstr *IGLPOptMI)
- : IGLPStrategy(DAG, TII, IGLPOptMI) {
+ MFMAExpSimpleInterleaveOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII)
+ : IGLPStrategy(DAG, TII) {
IsBottomUp = true;
}
};
@@ -2149,9 +2061,8 @@ class MFMASmallGemmSingleWaveOpt final : public IGLPStrategy {
return true;
}
- MFMASmallGemmSingleWaveOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII,
- const MachineInstr *IGLPOptMI)
- : IGLPStrategy(DAG, TII, IGLPOptMI) {
+ MFMASmallGemmSingleWaveOpt(ScheduleDAGInstrs *DAG, const SIInstrInfo *TII)
+ : IGLPStrategy(DAG, TII) {
IsBottomUp = false;
}
};
@@ -2166,22 +2077,6 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
unsigned DSWWithPermCount = 0;
unsigned DSWWithSharedVMEMCount = 0;
- bool IsInitial = Phase == AMDGPU::SchedulingPhase::Initial;
-
- if (!IsInitial) {
- const SIMachineFunctionInfo &MFI =
- *DAG->MF.getInfo<SIMachineFunctionInfo>();
- const SIMachineFunctionInfo::MFMASmallGemmSingleWaveCache *Cache =
- MFI.getMFMASmallGemmSingleWaveCache(IGLPOptMI);
- assert(Cache && "no cache found");
- DSWCount = Cache->DSWCount;
- DSWWithPermCount = Cache->DSWWithPermCount;
- DSWWithSharedVMEMCount = Cache->DSWWithSharedVMEMCount;
- }
-
- assert((!IsInitial || (DSWCount == 0 && DSWWithPermCount == 0 &&
- DSWWithSharedVMEMCount == 0)) &&
- "DSWCounters should be zero in pre-RA scheduling!");
SmallVector<SUnit *, 6> DSWithPerms;
for (auto &SU : DAG->SUnits) {
auto *I = SU.getInstr();
@@ -2190,7 +2085,7 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
else if (TII->isDS(*I)) {
if (I->mayLoad())
++DSRCount;
- else if (I->mayStore() && IsInitial) {
+ else if (I->mayStore()) {
++DSWCount;
for (auto Pred : SU.Preds) {
if (Pred.getSUnit()->getInstr()->getOpcode() ==
@@ -2203,62 +2098,55 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
}
}
- if (IsInitial) {
- DSWWithPermCount = DSWithPerms.size();
-
- // Get the count of DS_WRITES with V_PERM predecessors which
- // have loop carried dependencies (WAR) on the same VMEM_READs.
- // We consider partial overlap as a miss -- in other words,
- // for a given DS_W, we only consider another DS_W as matching
- // if there is a corresponding (in terms of the VMEM_R it uses) V_PERM pred
- // for every V_PERM pred of this DS_W.
- DenseMap<MachineInstr *, SUnit *> VMEMLookup;
- SmallVector<SUnit *, 6> Counted;
- for (SUnit *DSWrite : DSWithPerms) {
- SUnit *Cand = nullptr;
- bool MissedAny = false;
- for (auto &Pred : DSWrite->Preds) {
- if (Pred.getSUnit()->getInstr()->getOpcode() != AMDGPU::V_PERM_B32_e64)
- continue;
+ DSWWithPermCount = DSWithPerms.size();
+
+ // Get the count of DS_WRITES with V_PERM predecessors which
+ // have loop carried dependencies (WAR) on the same VMEM_READs.
+ // We consider partial overlap as a miss -- in other words,
+ // for a given DS_W, we only consider another DS_W as matching
+ // if there is a corresponding (in terms of the VMEM_R it uses) V_PERM pred
+ // for every V_PERM pred of this DS_W.
+ DenseMap<MachineInstr *, SUnit *> VMEMLookup;
+ SmallVector<SUnit *, 6> Counted;
+ for (SUnit *DSWrite : DSWithPerms) {
+ SUnit *Cand = nullptr;
+ bool MissedAny = false;
+ for (auto &Pred : DSWrite->Preds) {
+ if (Pred.getSUnit()->getInstr()->getOpcode() != AMDGPU::V_PERM_B32_e64)
+ continue;
- if (Cand && llvm::is_contained(Counted, Cand))
- break;
+ if (Cand && llvm::is_contained(Counted, Cand))
+ break;
- for (auto &Succ : Pred.getSUnit()->Succs) {
- auto *MI = Succ.getSUnit()->getInstr();
- if (!TII->isVMEM(*MI) || !MI->mayLoad())
- continue;
+ for (auto &Succ : Pred.getSUnit()->Succs) {
+ auto *MI = Succ.getSUnit()->getInstr();
+ if (!TII->isVMEM(*MI) || !MI->mayLoad())
+ continue;
- if (MissedAny || !VMEMLookup.size()) {
- MissedAny = true;
- VMEMLookup[MI] = DSWrite;
- continue;
- }
+ if (MissedAny || !VMEMLookup.size()) {
+ MissedAny = true;
+ VMEMLookup[MI] = DSWrite;
+ continue;
+ }
- auto [It, Inserted] = VMEMLookup.try_emplace(MI, DSWrite);
- if (Inserted) {
- MissedAny = true;
- continue;
- }
+ auto [It, Inserted] = VMEMLookup.try_emplace(MI, DSWrite);
+ if (Inserted) {
+ MissedAny = true;
+ continue;
+ }
- Cand = It->second;
- if (llvm::is_contained(Counted, Cand)) {
- MissedAny = true;
- break;
- }
+ Cand = It->second;
+ if (llvm::is_contained(Counted, Cand)) {
+ MissedAny = true;
+ break;
}
}
- if (!MissedAny && Cand) {
- DSWWithSharedVMEMCount += 2;
- Counted.push_back(Cand);
- Counted.push_back(DSWrite);
- }
}
-
- SIMachineFunctionInfo::MFMASmallGemmSingleWaveCache Cache = {
- DSWCount, DSWWithPermCount, DSWWithSharedVMEMCount};
- SIMachineFunctionInfo &MFI = *DAG->MF.getInfo<SIMachineFunctionInfo>();
- MFI.setMFMASmallGemmSingleWaveCache(IGLPOptMI, Cache);
+ if (!MissedAny && Cand) {
+ DSWWithSharedVMEMCount += 2;
+ Counted.push_back(Cand);
+ Counted.push_back(DSWrite);
+ }
}
assert(DSWWithSharedVMEMCount <= DSWWithPermCount);
@@ -2420,16 +2308,16 @@ bool MFMASmallGemmSingleWaveOpt::applyIGLPStrategy(
static std::unique_ptr<IGLPStrategy>
createIGLPStrategy(IGLPStrategyID ID, ScheduleDAGInstrs *DAG,
- const SIInstrInfo *TII, const MachineInstr *IGLPOptMI) {
+ const SIInstrInfo *TII) {
switch (ID) {
case MFMASmallGemmOptID:
- return std::make_unique<MFMASmallGemmOpt>(DAG, TII, IGLPOptMI);
+ return std::make_unique<MFMASmallGemmOpt>(DAG, TII);
case MFMASmallGemmSingleWaveOptID:
- return std::make_unique<MFMASmallGemmSingleWaveOpt>(DAG, TII, IGLPOptMI);
+ return std::make_unique<MFMASmallGemmSingleWaveOpt>(DAG, TII);
case MFMAExpInterleaveID:
- return std::make_unique<MFMAExpInterleaveOpt>(DAG, TII, IGLPOptMI);
+ return std::make_unique<MFMAExpInterleaveOpt>(DAG, TII);
case MFMAExpSimpleInterleaveID:
- return std::make_unique<MFMAExpSimpleInterleaveOpt>(DAG, TII, IGLPOptMI);
+ return std::make_unique<MFMAExpSimpleInterleaveOpt>(DAG, TII);
}
llvm_unreachable("Unknown IGLPStrategyID");
@@ -2807,7 +2695,7 @@ void IGroupLPDAGMutation::initSchedGroupBarrierPipelineStage(
bool IGroupLPDAGMutation::initIGLPOpt(SUnit &SU) {
IGLPStrategyID StrategyID =
(IGLPStrategyID)SU.getInstr()->getOperand(0).getImm();
- auto S = createIGLPStrategy(StrategyID, DAG, TII, SU.getInstr());
+ auto S = createIGLPStrategy(StrategyID, DAG, TII);
if (!S->shouldApplyStrategy(DAG, Phase))
return false;
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
index cb07f2da679c1..99fd55db33285 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.h
@@ -249,24 +249,6 @@ class GCNScheduleDAGMILive final : public ScheduleDAGMILive {
friend class ILPInitialScheduleStage;
friend class RegionPressureMap;
-public:
- struct IGLPExpInterleavePointerCache {
- /// MFMA instructions in the interleave pipeline that have no MFMA
- /// predecessors, i.e. the roots of independent MFMA chains.
- SmallVector<const MachineInstr *, 4> MFMAChainSeedInstrs;
- /// The first DS_READ that feeds into an MFMA chain seed.
- const MachineInstr *FirstPipeDSRInstr = nullptr;
- };
-
-private:
- /// Pass-local cache of pointer-based IGLP analysis data, keyed by the
- /// IGLP_OPT MachineInstr. This is safe because the DAG object lives for
- /// the duration of the pre-RA scheduling pass and the MachineInstrs are
- /// stable within that lifetime. Unlike SIMachineFunctionInfo caches, this
- /// is never serialized.
- DenseMap<const MachineInstr *, IGLPExpInterleavePointerCache>
- IGLPExpInterleavePtrCaches;
-
const GCNSubtarget &ST;
SIMachineFunctionInfo &MFI;
@@ -335,18 +317,6 @@ class GCNScheduleDAGMILive final : public ScheduleDAGMILive {
GCNScheduleDAGMILive(MachineSchedContext *C,
std::unique_ptr<MachineSchedStrategy> S);
- const IGLPExpInterleavePointerCache *
- getIGLPExpInterleavePointerCache(const MachineInstr *MI) const {
- auto It = IGLPExpInterleavePtrCaches.find(MI);
- return It == IGLPExpInterleavePtrCaches.end() ? nullptr : &It->second;
- }
-
- void
- setIGLPExpInterleavePointerCache(const MachineInstr *MI,
- const IGLPExpInterleavePointerCache &Cache) {
- IGLPExpInterleavePtrCaches[MI] = Cache;
- }
-
void schedule() override;
void finalizeSchedule() override;
diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
index 95867d442d278..32efae69b20c8 100644
--- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp
@@ -772,41 +772,6 @@ yaml::SIMachineFunctionInfo::SIMachineFunctionInfo(
auto SFI = MFI.getOptionalScavengeFI();
if (SFI)
ScavengeFI = yaml::FrameIndex(*SFI, MF.getFrameInfo());
-
- const DenseMap<const MachineInstr *,
- llvm::SIMachineFunctionInfo::MFMASmallGemmSingleWaveCache>
- &SmallGemmCaches = MFI.getMFMASmallGemmSingleWaveCaches();
- for (DenseMap<const MachineInstr *,
- llvm::SIMachineFunctionInfo::MFMASmallGemmSingleWaveCache>::
- const_iterator It = SmallGemmCaches.begin(),
- E = SmallGemmCaches.end();
- It != E; ++It) {
- const MachineInstr *MI = It->first;
- IGLPSmallGemmSingleWaveCaches.push_back(
- {static_cast<unsigned>(MI->getParent()->getNumber()),
- static_cast<unsigned>(MI->getOperand(0).getImm()), It->second.DSWCount,
- It->second.DSWWithPermCount, It->second.DSWWithSharedVMEMCount});
- }
-
- const DenseMap<const MachineInstr *,
- llvm::SIMachineFunctionInfo::MFMAExpInterleaveCache>
- &ExpInterleaveCaches = MFI.getMFMAExpInterleaveCaches();
- for (DenseMap<
- const MachineInstr *,
- llvm::SIMachineFunctionInfo::MFMAExpInterleaveCache>::const_iterator
- It = ExpInterleaveCaches.begin(),
- E = ExpInterleaveCaches.end();
- It != E; ++It) {
- const MachineInstr *MI = It->first;
- IGLPExpInterleaveCaches.push_back(
- {static_cast<unsigned>(MI->getParent()->getNumber()),
- static_cast<unsigned>(MI->getOperand(0).getImm()),
- It->second.TransPipeCount, It->second.MFMAPipeCount,
- It->second.AddPipeCount, It->second.MFMAEnablement,
- It->second.ExpRequirement, It->second.MFMAChains,
- It->second.MFMAChainLength, It->second.HasCvt,
- It->second.HasChainBetweenCvt, It->second.AnalysisResult});
- }
}
void yaml::SIMachineFunctionInfo::mappingImpl(yaml::IO &YamlIO) {
@@ -856,48 +821,6 @@ bool SIMachineFunctionInfo::initializeBaseYamlFields(
} else {
ScavengeFI = std::nullopt;
}
-
- if (!YamlMFI.IGLPSmallGemmSingleWaveCaches.empty() ||
- !YamlMFI.IGLPExpInterleaveCaches.empty()) {
- for (const MachineBasicBlock &MBB : MF) {
- for (const MachineInstr &MI : MBB) {
- if (MI.getOpcode() != AMDGPU::IGLP_OPT)
- continue;
- unsigned MBBNum = MBB.getNumber();
- unsigned StrategyId = static_cast<unsigned>(MI.getOperand(0).getImm());
-
- for (const yaml::IGLPSmallGemmSingleWaveCacheEntry &E :
- YamlMFI.IGLPSmallGemmSingleWaveCaches) {
- if (E.MBBNum == MBBNum && E.StrategyId == StrategyId) {
- MFMASmallGemmSingleWaveCache Cache = {
- E.DSWCount, E.DSWWithPermCount, E.DSWWithSharedVMEMCount};
- MFMASmallGemmSingleWaveCaches[&MI] = Cache;
- break;
- }
- }
-
- for (const yaml::IGLPExpInterleaveCacheEntry &E :
- YamlMFI.IGLPExpInterleaveCaches) {
- if (E.MBBNum == MBBNum && E.StrategyId == StrategyId) {
- MFMAExpInterleaveCache Cache;
- Cache.TransPipeCount = E.TransPipeCount;
- Cache.MFMAPipeCount = E.MFMAPipeCount;
- Cache.AddPipeCount = E.AddPipeCount;
- Cache.MFMAEnablement = E.MFMAEnablement;
- Cache.ExpRequirement = E.ExpRequirement;
- Cache.MFMAChains = E.MFMAChains;
- Cache.MFMAChainLength = E.MFMAChainLength;
- Cache.HasCvt = E.HasCvt;
- Cache.HasChainBetweenCvt = E.HasChainBetweenCvt;
- Cache.AnalysisResult = E.AnalysisResult;
- MFMAExpInterleaveCaches[&MI] = Cache;
- break;
- }
- }
- }
- }
- }
-
return false;
}
diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
index 7a2ac80c002d1..617862db8f506 100644
--- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
@@ -20,7 +20,6 @@
#include "MCTargetDesc/AMDGPUMCTargetDesc.h"
#include "SIInstrInfo.h"
#include "SIModeRegisterDefaults.h"
-#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/SetVector.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/CodeGen/MIRYamlMapping.h"
@@ -30,8 +29,6 @@
namespace llvm {
-class MachineInstr;
-
class MachineFrameInfo;
class MachineFunction;
class SIMachineFunctionInfo;
@@ -262,56 +259,6 @@ template <> struct MappingTraits<SIMode> {
}
};
-struct IGLPSmallGemmSingleWaveCacheEntry {
- unsigned MBBNum = 0;
- unsigned StrategyId = 0;
- unsigned DSWCount = 0;
- unsigned DSWWithPermCount = 0;
- unsigned DSWWithSharedVMEMCount = 0;
-};
-
-template <> struct MappingTraits<IGLPSmallGemmSingleWaveCacheEntry> {
- static void mapping(IO &YamlIO, IGLPSmallGemmSingleWaveCacheEntry &E) {
- YamlIO.mapRequired("mbb", E.MBBNum);
- YamlIO.mapRequired("strategyId", E.StrategyId);
- YamlIO.mapRequired("dswCount", E.DSWCount);
- YamlIO.mapRequired("dswWithPermCount", E.DSWWithPermCount);
- YamlIO.mapRequired("dswWithSharedVMEMCount", E.DSWWithSharedVMEMCount);
- }
-};
-
-struct IGLPExpInterleaveCacheEntry {
- unsigned MBBNum = 0;
- unsigned StrategyId = 0;
- unsigned TransPipeCount = 0;
- unsigned MFMAPipeCount = 0;
- unsigned AddPipeCount = 0;
- unsigned MFMAEnablement = 0;
- unsigned ExpRequirement = 0;
- unsigned MFMAChains = 0;
- unsigned MFMAChainLength = 0;
- bool HasCvt = false;
- bool HasChainBetweenCvt = false;
- bool AnalysisResult = false;
-};
-
-template <> struct MappingTraits<IGLPExpInterleaveCacheEntry> {
- static void mapping(IO &YamlIO, IGLPExpInterleaveCacheEntry &E) {
- YamlIO.mapRequired("mbb", E.MBBNum);
- YamlIO.mapRequired("strategyId", E.StrategyId);
- YamlIO.mapRequired("transPipeCount", E.TransPipeCount);
- YamlIO.mapRequired("mfmaPipeCount", E.MFMAPipeCount);
- YamlIO.mapRequired("addPipeCount", E.AddPipeCount);
- YamlIO.mapRequired("mfmaEnablement", E.MFMAEnablement);
- YamlIO.mapRequired("expRequirement", E.ExpRequirement);
- YamlIO.mapRequired("mfmaChains", E.MFMAChains);
- YamlIO.mapRequired("mfmaChainLength", E.MFMAChainLength);
- YamlIO.mapRequired("hasCvt", E.HasCvt);
- YamlIO.mapRequired("hasChainBetweenCvt", E.HasChainBetweenCvt);
- YamlIO.mapRequired("analysisResult", E.AnalysisResult);
- }
-};
-
struct SIMachineFunctionInfo final : public yaml::MachineFunctionInfo {
uint64_t ExplicitKernArgSize = 0;
Align MaxKernArgAlign;
@@ -356,9 +303,6 @@ struct SIMachineFunctionInfo final : public yaml::MachineFunctionInfo {
bool HasInitWholeWave = false;
bool IsWholeWaveFunction = false;
- SmallVector<IGLPSmallGemmSingleWaveCacheEntry> IGLPSmallGemmSingleWaveCaches;
- SmallVector<IGLPExpInterleaveCacheEntry> IGLPExpInterleaveCaches;
-
unsigned DynamicVGPRBlockSize = 0;
unsigned ScratchReservedForDynamicVGPRs = 0;
@@ -416,9 +360,6 @@ template <> struct MappingTraits<SIMachineFunctionInfo> {
YamlIO.mapOptional("longBranchReservedReg", MFI.LongBranchReservedReg,
StringValue());
YamlIO.mapOptional("hasInitWholeWave", MFI.HasInitWholeWave, false);
- YamlIO.mapOptional("iglpSmallGemmSingleWaveCaches",
- MFI.IGLPSmallGemmSingleWaveCaches);
- YamlIO.mapOptional("iglpExpInterleaveCaches", MFI.IGLPExpInterleaveCaches);
YamlIO.mapOptional("dynamicVGPRBlockSize", MFI.DynamicVGPRBlockSize, false);
YamlIO.mapOptional("scratchReservedForDynamicVGPRs",
MFI.ScratchReservedForDynamicVGPRs, 0);
@@ -609,25 +550,6 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
bool IsDead = false;
};
- struct MFMASmallGemmSingleWaveCache {
- unsigned DSWCount = 0;
- unsigned DSWWithPermCount = 0;
- unsigned DSWWithSharedVMEMCount = 0;
- };
-
- struct MFMAExpInterleaveCache {
- unsigned TransPipeCount = 0;
- unsigned MFMAPipeCount = 0;
- unsigned AddPipeCount = 0;
- unsigned MFMAEnablement = 0;
- unsigned ExpRequirement = 0;
- unsigned MFMAChains = 0;
- unsigned MFMAChainLength = 0;
- bool HasCvt = false;
- bool HasChainBetweenCvt = false;
- bool AnalysisResult = false;
- };
-
private:
// To track virtual VGPR + lane index for each subregister of the SGPR spilled
// to frameindex key during SILowerSGPRSpills pass.
@@ -692,11 +614,7 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
// load/store is enabled.
IndexedMap<uint32_t, VGPRBlock2IndexFunctor> MaskForVGPRBlockOps;
- DenseMap<const MachineInstr *, MFMASmallGemmSingleWaveCache>
- MFMASmallGemmSingleWaveCaches;
- DenseMap<const MachineInstr *, MFMAExpInterleaveCache>
- MFMAExpInterleaveCaches;
-
+private:
Register VGPRForAGPRCopy;
bool allocateVirtualVGPRForSGPRSpills(MachineFunction &MF, int FI,
@@ -714,41 +632,6 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
VGPRForAGPRCopy = NewVGPRForAGPRCopy;
}
- const MFMASmallGemmSingleWaveCache *
- getMFMASmallGemmSingleWaveCache(const MachineInstr *MI) const {
- DenseMap<const MachineInstr *, MFMASmallGemmSingleWaveCache>::const_iterator
- It = MFMASmallGemmSingleWaveCaches.find(MI);
- return It == MFMASmallGemmSingleWaveCaches.end() ? nullptr : &It->second;
- }
-
- void
- setMFMASmallGemmSingleWaveCache(const MachineInstr *MI,
- const MFMASmallGemmSingleWaveCache &Cache) {
- MFMASmallGemmSingleWaveCaches[MI] = Cache;
- }
-
- const MFMAExpInterleaveCache *
- getMFMAExpInterleaveCache(const MachineInstr *MI) const {
- DenseMap<const MachineInstr *, MFMAExpInterleaveCache>::const_iterator It =
- MFMAExpInterleaveCaches.find(MI);
- return It == MFMAExpInterleaveCaches.end() ? nullptr : &It->second;
- }
-
- void setMFMAExpInterleaveCache(const MachineInstr *MI,
- const MFMAExpInterleaveCache &Cache) {
- MFMAExpInterleaveCaches[MI] = Cache;
- }
-
- const DenseMap<const MachineInstr *, MFMASmallGemmSingleWaveCache> &
- getMFMASmallGemmSingleWaveCaches() const {
- return MFMASmallGemmSingleWaveCaches;
- }
-
- const DenseMap<const MachineInstr *, MFMAExpInterleaveCache> &
- getMFMAExpInterleaveCaches() const {
- return MFMAExpInterleaveCaches;
- }
-
bool isCalleeSavedReg(const MCPhysReg *CSRegs, MCPhysReg Reg) const;
void setMaskForVGPRBlockOps(Register RegisterBlock, uint32_t Mask) {
@@ -764,6 +647,7 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
return MaskForVGPRBlockOps.inBounds(RegisterBlock);
}
+public:
SIMachineFunctionInfo(const SIMachineFunctionInfo &MFI) = default;
SIMachineFunctionInfo(const Function &F, const GCNSubtarget *STI);
@@ -1345,7 +1229,4 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction,
} // end namespace llvm
-LLVM_YAML_IS_SEQUENCE_VECTOR(llvm::yaml::IGLPSmallGemmSingleWaveCacheEntry)
-LLVM_YAML_IS_SEQUENCE_VECTOR(llvm::yaml::IGLPExpInterleaveCacheEntry)
-
#endif // LLVM_LIB_TARGET_AMDGPU_SIMACHINEFUNCTIONINFO_H
diff --git a/llvm/test/CodeGen/AMDGPU/iglp-cache-serialization.mir b/llvm/test/CodeGen/AMDGPU/iglp-cache-serialization.mir
deleted file mode 100644
index 501eb308ffc48..0000000000000
--- a/llvm/test/CodeGen/AMDGPU/iglp-cache-serialization.mir
+++ /dev/null
@@ -1,85 +0,0 @@
-# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -run-pass=none -o - %s | FileCheck -check-prefix=ROUNDTRIP %s
-# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -run-pass=postmisched -o - %s | FileCheck -check-prefix=POSTRA %s
-
-# Test that IGLP cache data serializes/deserializes through MIR and is
-# available for PostRA scheduling.
-
-# Verify round-trip of the serialized cache fields.
-# ROUNDTRIP: iglpSmallGemmSingleWaveCaches:
-# ROUNDTRIP-NEXT: - mbb: 1
-# ROUNDTRIP-NEXT: strategyId: 1
-# ROUNDTRIP-NEXT: dswCount: 6
-# ROUNDTRIP-NEXT: dswWithPermCount: 2
-# ROUNDTRIP-NEXT: dswWithSharedVMEMCount: 0
-# ROUNDTRIP: iglpExpInterleaveCaches:
-# ROUNDTRIP-NEXT: - mbb: 0
-# ROUNDTRIP-NEXT: strategyId: 2
-# ROUNDTRIP-NEXT: transPipeCount: 32
-# ROUNDTRIP-NEXT: mfmaPipeCount: 40
-# ROUNDTRIP-NEXT: addPipeCount: 33
-# ROUNDTRIP-NEXT: mfmaEnablement: 2
-# ROUNDTRIP-NEXT: expRequirement: 4
-# ROUNDTRIP-NEXT: mfmaChains: 2
-# ROUNDTRIP-NEXT: mfmaChainLength: 20
-# ROUNDTRIP-NEXT: hasCvt: true
-# ROUNDTRIP-NEXT: hasChainBetweenCvt: false
-# ROUNDTRIP-NEXT: analysisResult: true
-
-# Verify PostRA scheduling doesn't crash and passes through the IGLP_OPTs.
-# POSTRA: IGLP_OPT 2
-# POSTRA: V_MFMA
-# POSTRA: IGLP_OPT 1
---- |
- define amdgpu_kernel void @iglp_cache_test() #0 { ret void }
-
- attributes #0 = { "amdgpu-flat-work-group-size"="256,256" }
-...
----
-name: iglp_cache_test
-alignment: 1
-tracksRegLiveness: true
-noPhis: true
-isSSA: false
-noVRegs: true
-machineFunctionInfo:
- isEntryFunction: false
- occupancy: 5
- iglpExpInterleaveCaches:
- - mbb: 0
- strategyId: 2
- transPipeCount: 32
- mfmaPipeCount: 40
- addPipeCount: 33
- mfmaEnablement: 2
- expRequirement: 4
- mfmaChains: 2
- mfmaChainLength: 20
- hasCvt: true
- hasChainBetweenCvt: false
- analysisResult: true
- iglpSmallGemmSingleWaveCaches:
- - mbb: 1
- strategyId: 1
- dswCount: 6
- dswWithPermCount: 2
- dswWithSharedVMEMCount: 0
-body: |
- bb.0:
- liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr36_vgpr37, $vgpr40_vgpr41, $vgpr44_vgpr45, $vgpr38_vgpr39, $vgpr42_vgpr43
-
- IGLP_OPT 2
- $vgpr50 = V_EXP_F32_e32 $vgpr0, implicit $mode, implicit $exec
- $vgpr51 = V_EXP_F32_e32 $vgpr1, implicit $mode, implicit $exec
- $vgpr52 = V_ADD_F32_e32 $vgpr2, $vgpr3, implicit $mode, implicit $exec
- early-clobber $vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23_vgpr24_vgpr25_vgpr26_vgpr27_vgpr28_vgpr29_vgpr30_vgpr31 = contract V_MFMA_F32_32X32X8F16_vgprcd_e64 $vgpr36_vgpr37, $vgpr40_vgpr41, 0, 0, 0, 0, implicit $mode, implicit $exec
- early-clobber $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15 = contract V_MFMA_F32_32X32X8F16_vgprcd_e64 $vgpr44_vgpr45, $vgpr40_vgpr41, 0, 0, 0, 0, implicit $mode, implicit $exec
- S_BRANCH %bb.1
-
- bb.1:
- liveins: $vgpr0, $vgpr1
-
- IGLP_OPT 1
- $vgpr0 = V_ADD_F32_e32 $vgpr0, $vgpr1, implicit $mode, implicit $exec
- S_ENDPGM 0
-
-...
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.AFLCustomIRMutator.opt.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.AFLCustomIRMutator.opt.ll
index 9dad8ff365651..4005aacc90ef3 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.AFLCustomIRMutator.opt.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.AFLCustomIRMutator.opt.ll
@@ -5,21 +5,16 @@ define amdgpu_kernel void @test_iglp_opt_rev_mfma_gemm(<1 x i64> %L1) {
; GCN-LABEL: test_iglp_opt_rev_mfma_gemm:
; GCN: ; %bb.0: ; %entry
; GCN-NEXT: v_mov_b32_e32 v0, 0
-; GCN-NEXT: ds_read_b128 v[2:5], v0
-; GCN-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0
; GCN-NEXT: ds_read_b128 v[30:33], v0 offset:112
; GCN-NEXT: ds_read_b128 v[26:29], v0 offset:96
; GCN-NEXT: ds_read_b128 v[22:25], v0 offset:80
; GCN-NEXT: ds_read_b128 v[18:21], v0 offset:64
+; GCN-NEXT: ds_read_b128 v[2:5], v0
; GCN-NEXT: ds_read_b128 v[6:9], v0 offset:16
; GCN-NEXT: ds_read_b128 v[10:13], v0 offset:32
; GCN-NEXT: ds_read_b128 v[14:17], v0 offset:48
+; GCN-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0
; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: ds_write_b128 v0, v[2:5]
-; GCN-NEXT: v_mov_b32_e32 v2, 0
-; GCN-NEXT: v_mov_b32_e32 v3, 0
-; GCN-NEXT: s_cmp_lg_u64 s[0:1], 0
-; GCN-NEXT: ; iglp_opt mask(0x00000001)
; GCN-NEXT: ds_write_b128 v0, v[30:33] offset:112
; GCN-NEXT: ds_write_b128 v0, v[26:29] offset:96
; GCN-NEXT: ds_write_b128 v0, v[22:25] offset:80
@@ -27,6 +22,11 @@ define amdgpu_kernel void @test_iglp_opt_rev_mfma_gemm(<1 x i64> %L1) {
; GCN-NEXT: ds_write_b128 v0, v[14:17] offset:48
; GCN-NEXT: ds_write_b128 v0, v[10:13] offset:32
; GCN-NEXT: ds_write_b128 v0, v[6:9] offset:16
+; GCN-NEXT: ds_write_b128 v0, v[2:5]
+; GCN-NEXT: v_mov_b32_e32 v2, 0
+; GCN-NEXT: v_mov_b32_e32 v3, 0
+; GCN-NEXT: s_cmp_lg_u64 s[0:1], 0
+; GCN-NEXT: ; iglp_opt mask(0x00000001)
; GCN-NEXT: ds_write_b64 v0, v[2:3]
; GCN-NEXT: s_endpgm
entry:
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 94de6dd31cad5..729596f541147 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
@@ -79,7 +79,6 @@
; 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
@@ -87,6 +86,7 @@
; 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: $vgpr0_vgpr1_vgpr2_vgpr3_vgpr4_vgpr5_vgpr6_vgpr7_vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15
; GCN-NEXT: ; implicit-def: $vgpr21
; GCN-NEXT: ; implicit-def: $vgpr22
; GCN-NEXT: ; implicit-def: $vgpr23
@@ -510,69 +510,70 @@
; GCN-NEXT: v_fma_f32 v113, s4, v116, -v128
; GCN-NEXT: v_mul_f32_e32 v141, 0x3fb8aa3b, v113
; GCN-NEXT: v_fma_f32 v113, s4, v117, -v128
+ ; GCN-NEXT: v_fma_f32 v112, s4, v112, -v128
; GCN-NEXT: v_mul_f32_e32 v142, 0x3fb8aa3b, v113
; GCN-NEXT: v_fma_f32 v113, s4, v118, -v128
- ; GCN-NEXT: v_fma_f32 v112, s4, v112, -v128
+ ; GCN-NEXT: v_mul_f32_e32 v112, 0x3fb8aa3b, v112
; GCN-NEXT: v_mul_f32_e32 v143, 0x3fb8aa3b, v113
; GCN-NEXT: v_fma_f32 v113, s4, v119, -v128
- ; GCN-NEXT: v_fma_f32 v118, s4, v120, -v128
- ; GCN-NEXT: v_fma_f32 v120, s4, v121, -v128
- ; GCN-NEXT: v_mul_f32_e32 v112, 0x3fb8aa3b, v112
; GCN-NEXT: v_mul_f32_e32 v144, 0x3fb8aa3b, v113
- ; GCN-NEXT: v_mul_f32_e32 v149, 0x3fb8aa3b, v120
- ; GCN-NEXT: v_fma_f32 v120, s4, v122, -v128
+ ; GCN-NEXT: v_exp_f32_e32 v113, v112
; GCN-NEXT: v_exp_f32_e32 v114, v138
; GCN-NEXT: v_exp_f32_e32 v115, v139
; GCN-NEXT: v_exp_f32_e32 v116, v140
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v112, v113
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v119, v114
+ ; GCN-NEXT: v_fma_f32 v118, s4, v120, -v128
+ ; GCN-NEXT: v_fma_f32 v120, s4, v121, -v128
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v121, v116
+ ; GCN-NEXT: v_pack_b32_f16 v146, v112, v119
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v112, v115
+ ; GCN-NEXT: v_sub_f32_e32 v129, v211, v128
+ ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v129
+ ; GCN-NEXT: v_mul_f32_e32 v149, 0x3fb8aa3b, v120
+ ; GCN-NEXT: v_fma_f32 v120, s4, v122, -v128
; GCN-NEXT: v_exp_f32_e32 v117, v141
; GCN-NEXT: v_mul_f32_e32 v148, 0x3fb8aa3b, v118
; GCN-NEXT: v_exp_f32_e32 v118, v142
+ ; GCN-NEXT: v_exp_f32_e32 v119, v143
; GCN-NEXT: v_mul_f32_e32 v150, 0x3fb8aa3b, v120
; GCN-NEXT: v_exp_f32_e32 v120, v144
- ; GCN-NEXT: v_exp_f32_e32 v113, v112
- ; GCN-NEXT: v_cvt_f16_f32_e32 v119, v114
- ; GCN-NEXT: v_cvt_f16_f32_e32 v121, v116
- ; GCN-NEXT: v_sub_f32_e32 v129, v211, v128
- ; GCN-NEXT: v_cvt_f16_f32_e32 v112, v113
- ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v129
+ ; GCN-NEXT: v_pack_b32_f16 v147, v112, v121
+ ; GCN-NEXT: v_exp_f32_e32 v112, v129
; GCN-NEXT: ds_read_b128 v[138:141], v198 offset:1152
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: ds_read_b128 v[142:145], v198 offset:1728
+ ; GCN-NEXT: s_waitcnt lgkmcnt(0)
+ ; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: v_fma_f32 v122, s4, v123, -v128
- ; GCN-NEXT: v_pack_b32_f16 v146, v112, v119
- ; GCN-NEXT: v_cvt_f16_f32_e32 v112, v115
; GCN-NEXT: v_mul_f32_e32 v151, 0x3fb8aa3b, v122
- ; GCN-NEXT: v_cvt_f16_f32_e32 v123, v117
- ; GCN-NEXT: v_fma_f32 v122, s4, v124, -v128
- ; GCN-NEXT: v_pack_b32_f16 v147, v112, v121
- ; GCN-NEXT: v_exp_f32_e32 v112, v129
- ; GCN-NEXT: v_cvt_f16_f32_e32 v124, v118
- ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v122
- ; GCN-NEXT: v_fma_f32 v125, s4, v125, -v128
- ; GCN-NEXT: v_pk_mul_f32 v[0:1], v[0:1], v[112:113] op_sel_hi:[1,0]
- ; GCN-NEXT: v_pk_mul_f32 v[2:3], v[2:3], v[112:113] op_sel_hi:[1,0]
- ; GCN-NEXT: v_pk_mul_f32 v[4:5], v[4:5], v[112:113] op_sel_hi:[1,0]
- ; GCN-NEXT: v_pk_mul_f32 v[6:7], v[6:7], v[112:113] op_sel_hi:[1,0]
- ; GCN-NEXT: v_pk_mul_f32 v[8:9], v[8:9], v[112:113] op_sel_hi:[1,0]
- ; GCN-NEXT: v_pk_mul_f32 v[10:11], v[10:11], v[112:113] op_sel_hi:[1,0]
- ; GCN-NEXT: v_pk_mul_f32 v[12:13], v[12:13], v[112:113] op_sel_hi:[1,0]
- ; GCN-NEXT: v_pk_mul_f32 v[14:15], v[14:15], v[112:113] op_sel_hi:[1,0]
; GCN-NEXT: v_pk_mul_f32 v[32:33], v[32:33], v[112:113] op_sel_hi:[1,0]
; GCN-NEXT: v_pk_mul_f32 v[34:35], v[34:35], v[112:113] op_sel_hi:[1,0]
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[130:131], v[146:147], v[0:15]
- ; GCN-NEXT: v_exp_f32_e32 v119, v143
- ; GCN-NEXT: ds_read_b128 v[142:145], v198 offset:1728
- ; GCN-NEXT: s_waitcnt lgkmcnt(0)
- ; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: v_pk_mul_f32 v[36:37], v[36:37], v[112:113] op_sel_hi:[1,0]
; GCN-NEXT: v_pk_mul_f32 v[38:39], v[38:39], v[112:113] op_sel_hi:[1,0]
; GCN-NEXT: v_pk_mul_f32 v[40:41], v[40:41], v[112:113] op_sel_hi:[1,0]
; GCN-NEXT: v_pk_mul_f32 v[42:43], v[42:43], v[112:113] op_sel_hi:[1,0]
; GCN-NEXT: v_pk_mul_f32 v[44:45], v[44:45], v[112:113] op_sel_hi:[1,0]
; GCN-NEXT: v_pk_mul_f32 v[46:47], v[46:47], v[112:113] op_sel_hi:[1,0]
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v123, v117
+ ; GCN-NEXT: v_fma_f32 v122, s4, v124, -v128
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v124, v118
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[134:135], v[146:147], v[32:47]
+ ; GCN-NEXT: v_mul_f32_e64 v0, v0, v112
+ ; GCN-NEXT: v_mul_f32_e64 v1, v1, v112
+ ; GCN-NEXT: v_mul_f32_e64 v2, v2, v112
+ ; GCN-NEXT: v_mul_f32_e64 v3, v3, v112
+ ; GCN-NEXT: v_mul_f32_e64 v4, v4, v112
+ ; GCN-NEXT: v_mul_f32_e64 v5, v5, v112
+ ; GCN-NEXT: v_pk_mul_f32 v[6:7], v[6:7], v[112:113] op_sel_hi:[1,0]
+ ; GCN-NEXT: v_pk_mul_f32 v[8:9], v[8:9], v[112:113] op_sel_hi:[1,0]
+ ; GCN-NEXT: v_pk_mul_f32 v[10:11], v[10:11], v[112:113] op_sel_hi:[1,0]
+ ; GCN-NEXT: v_pk_mul_f32 v[12:13], v[12:13], v[112:113] op_sel_hi:[1,0]
+ ; GCN-NEXT: v_pk_mul_f32 v[14:15], v[14:15], v[112:113] op_sel_hi:[1,0]
; GCN-NEXT: v_pk_mul_f32 v[16:17], v[16:17], v[112:113] op_sel_hi:[1,0]
; GCN-NEXT: v_pk_mul_f32 v[18:19], v[18:19], v[112:113] op_sel_hi:[1,0]
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[134:135], v[146:147], v[32:47]
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[130:131], v[146:147], v[0:15]
; GCN-NEXT: v_mul_f32_e64 v20, v20, v112
; GCN-NEXT: v_mul_f32_e64 v21, v21, v112
; GCN-NEXT: v_mul_f32_e64 v22, v22, v112
@@ -595,46 +596,50 @@
; GCN-NEXT: v_fma_f32 v124, s4, v126, -v128
; GCN-NEXT: v_cvt_f16_f32_e32 v126, v120
; GCN-NEXT: v_exp_f32_e32 v121, v148
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[138:139], v[146:147], v[16:31]
+ ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v122
; GCN-NEXT: v_exp_f32_e32 v122, v149
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[138:139], v[146:147], v[16:31]
; GCN-NEXT: v_pack_b32_f16 v135, v130, v126
+ ; GCN-NEXT: v_exp_f32_e32 v123, v150
; GCN-NEXT: v_mul_f32_e32 v138, 0x3fb8aa3b, v124
+ ; GCN-NEXT: v_exp_f32_e32 v124, v151
; GCN-NEXT: v_cvt_f16_f32_e32 v126, v121
+ ; GCN-NEXT: v_fma_f32 v125, s4, v125, -v128
; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v125
- ; GCN-NEXT: v_fma_f32 v139, s4, v96, -v128
- ; GCN-NEXT: v_fma_f32 v127, s4, v127, -v128
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[142:143], v[146:147], v[48:63]
- ; GCN-NEXT: v_exp_f32_e32 v123, v150
+ ; GCN-NEXT: v_fma_f32 v127, s4, v127, -v128
; GCN-NEXT: v_mul_f32_e32 v127, 0x3fb8aa3b, v127
+ ; GCN-NEXT: v_fma_f32 v139, s4, v96, -v128
+ ; GCN-NEXT: v_exp_f32_e32 v96, v129
+ ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v139
; GCN-NEXT: v_fma_f32 v143, s4, v101, -v128
; GCN-NEXT: v_fma_f32 v64, s4, v64, -v128
- ; GCN-NEXT: v_fma_f32 v65, s4, v65, -v128
- ; GCN-NEXT: v_fma_f32 v68, s4, v68, -v128
- ; GCN-NEXT: v_fma_f32 v69, s4, v69, -v128
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[136:137], v[134:135], v[32:47]
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v136, v122
+ ; GCN-NEXT: v_fma_f32 v137, s4, v97, -v128
+ ; GCN-NEXT: v_exp_f32_e32 v97, v125
+ ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v137
+ ; GCN-NEXT: v_pack_b32_f16 v126, v126, v136
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v136, v123
+ ; GCN-NEXT: v_fma_f32 v137, s4, v98, -v128
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[134:135], v[0:15]
- ; GCN-NEXT: v_exp_f32_e32 v124, v151
; GCN-NEXT: ds_read_b128 v[130:133], v197
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: ds_read_b128 v[146:149], v197 offset:576
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[136:137], v[134:135], v[32:47]
- ; GCN-NEXT: v_cvt_f16_f32_e32 v136, v122
- ; GCN-NEXT: v_exp_f32_e32 v96, v129
- ; GCN-NEXT: v_fma_f32 v137, s4, v97, -v128
- ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v139
- ; GCN-NEXT: v_pack_b32_f16 v126, v126, v136
- ; GCN-NEXT: v_cvt_f16_f32_e32 v136, v123
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[140:141], v[134:135], v[16:31]
- ; GCN-NEXT: v_exp_f32_e32 v97, v125
- ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v137
- ; GCN-NEXT: v_fma_f32 v137, s4, v98, -v128
; GCN-NEXT: v_mul_f32_e32 v142, 0x3fb8aa3b, v137
+ ; GCN-NEXT: v_exp_f32_e32 v98, v138
+ ; GCN-NEXT: v_exp_f32_e32 v101, v125
+ ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v143
+ ; GCN-NEXT: v_fma_f32 v65, s4, v65, -v128
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[140:141], v[134:135], v[16:31]
+ ; GCN-NEXT: v_fma_f32 v68, s4, v68, -v128
+ ; GCN-NEXT: v_fma_f32 v69, s4, v69, -v128
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[144:145], v[134:135], v[48:63]
; GCN-NEXT: v_cvt_f16_f32_e32 v134, v124
; GCN-NEXT: v_fma_f32 v135, s4, v99, -v128
- ; GCN-NEXT: v_exp_f32_e32 v98, v138
; GCN-NEXT: v_exp_f32_e32 v99, v127
; GCN-NEXT: v_mul_f32_e32 v150, 0x3fb8aa3b, v135
; GCN-NEXT: v_pack_b32_f16 v127, v136, v134
@@ -665,28 +670,26 @@
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GCN-NEXT: ds_write_b64 v202, v[194:195]
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[134:135], v[126:127], v[16:31]
+ ; GCN-NEXT: v_fma_f32 v134, s4, v102, -v128
+ ; GCN-NEXT: v_mul_f32_e32 v156, 0x3fb8aa3b, v134
+ ; GCN-NEXT: v_exp_f32_e32 v102, v142
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[146:147], v[126:127], v[32:47]
- ; GCN-NEXT: v_exp_f32_e32 v101, v125
; GCN-NEXT: v_pack_b32_f16 v146, v130, v131
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_load_dwordx2 v[130:131], v210, s[0:3], 0 offen sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v143
- ; GCN-NEXT: v_cvt_f16_f32_e32 v147, v98
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[134:135], v[126:127], v[16:31]
- ; GCN-NEXT: v_fma_f32 v134, s4, v102, -v128
- ; GCN-NEXT: v_mul_f32_e32 v156, 0x3fb8aa3b, v134
; GCN-NEXT: buffer_load_dwordx2 v[134:135], v207, s[0:3], 0 offen sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_exp_f32_e32 v102, v142
; GCN-NEXT: buffer_load_dwordx2 v[142:143], v208, s[0:3], 0 offen sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: buffer_load_dwordx2 v[144:145], v209, s[0:3], 0 offen sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v147, v98
; GCN-NEXT: ;;#ASMSTART
; GCN-NEXT: s_waitcnt vmcnt(8)
; GCN-NEXT: ;;#ASMEND
@@ -694,47 +697,47 @@
; GCN-NEXT: v_cvt_f16_f32_e32 v126, v99
; GCN-NEXT: v_fma_f32 v127, s4, v103, -v128
; GCN-NEXT: v_exp_f32_e32 v103, v150
- ; GCN-NEXT: v_fma_f32 v139, s4, v105, -v128
- ; GCN-NEXT: v_pack_b32_f16 v147, v147, v126
; GCN-NEXT: v_mul_f32_e32 v138, 0x3fb8aa3b, v127
+ ; GCN-NEXT: v_pack_b32_f16 v147, v147, v126
+ ; GCN-NEXT: v_fma_f32 v139, s4, v105, -v128
+ ; GCN-NEXT: v_exp_f32_e32 v105, v125
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[148:149], v[146:147], v[32:47]
+ ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v139
; GCN-NEXT: v_perm_b32 v152, v135, v131, s5
; GCN-NEXT: v_perm_b32 v154, v135, v131, s7
; GCN-NEXT: v_fma_f32 v135, s4, v104, -v128
; GCN-NEXT: v_perm_b32 v126, v134, v130, s5
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[146:147], v[0:15]
; GCN-NEXT: v_perm_b32 v150, v134, v130, s7
; GCN-NEXT: v_cvt_f16_f32_e32 v134, v100
; GCN-NEXT: v_exp_f32_e32 v104, v129
; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v135
; GCN-NEXT: v_cvt_f16_f32_e32 v135, v101
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[146:147], v[0:15]
; GCN-NEXT: ds_read_b128 v[130:133], v198
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: v_perm_b32 v127, v144, v142, s5
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[148:149], v[146:147], v[32:47]
; GCN-NEXT: v_pack_b32_f16 v148, v134, v135
; GCN-NEXT: v_fma_f32 v135, s4, v106, -v128
- ; GCN-NEXT: v_exp_f32_e32 v105, v125
; GCN-NEXT: v_cvt_f16_f32_e32 v134, v102
- ; GCN-NEXT: v_perm_b32 v151, v144, v142, s7
- ; GCN-NEXT: v_perm_b32 v153, v145, v143, s5
- ; GCN-NEXT: v_perm_b32 v155, v145, v143, s7
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[136:137], v[146:147], v[16:31]
; GCN-NEXT: v_exp_f32_e32 v106, v156
; GCN-NEXT: v_mul_f32_e32 v156, 0x3fb8aa3b, v135
; GCN-NEXT: v_cvt_f16_f32_e32 v135, v103
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[136:137], v[146:147], v[16:31]
; GCN-NEXT: v_fma_f32 v136, s4, v107, -v128
+ ; GCN-NEXT: v_perm_b32 v151, v144, v142, s7
+ ; GCN-NEXT: v_perm_b32 v153, v145, v143, s5
+ ; GCN-NEXT: v_perm_b32 v155, v145, v143, s7
; GCN-NEXT: ds_read_b128 v[142:145], v198 offset:576
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v139
; GCN-NEXT: v_pack_b32_f16 v149, v134, v135
+ ; GCN-NEXT: v_exp_f32_e32 v107, v138
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[140:141], v[146:147], v[48:63]
; GCN-NEXT: v_mul_f32_e32 v146, 0x3fb8aa3b, v136
; GCN-NEXT: ds_read_b128 v[134:137], v198 offset:1152
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_exp_f32_e32 v107, v138
; GCN-NEXT: ds_read_b128 v[138:141], v198 offset:1728
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
@@ -744,6 +747,10 @@
; GCN-NEXT: v_exp_f32_e32 v108, v129
; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v131
; GCN-NEXT: v_cvt_f16_f32_e32 v131, v105
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[134:135], v[148:149], v[16:31]
+ ; GCN-NEXT: v_fma_f32 v135, s4, v111, -v128
+ ; GCN-NEXT: v_exp_f32_e32 v111, v146
+ ; GCN-NEXT: v_mul_f32_e32 v135, 0x3fb8aa3b, v135
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[142:143], v[148:149], v[32:47]
; GCN-NEXT: v_fma_f32 v142, s4, v109, -v128
; GCN-NEXT: v_exp_f32_e32 v109, v125
@@ -751,32 +758,16 @@
; GCN-NEXT: v_pack_b32_f16 v142, v130, v131
; GCN-NEXT: v_fma_f32 v131, s4, v110, -v128
; GCN-NEXT: v_cvt_f16_f32_e32 v130, v106
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[134:135], v[148:149], v[16:31]
; GCN-NEXT: v_mul_f32_e32 v134, 0x3fb8aa3b, v131
; GCN-NEXT: v_cvt_f16_f32_e32 v131, v107
; GCN-NEXT: v_exp_f32_e32 v110, v156
- ; GCN-NEXT: v_fma_f32 v135, s4, v111, -v128
- ; GCN-NEXT: v_mul_f32_e32 v135, 0x3fb8aa3b, v135
- ; GCN-NEXT: v_pack_b32_f16 v143, v130, v131
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[138:139], v[148:149], v[48:63]
- ; GCN-NEXT: v_exp_f32_e32 v111, v146
; GCN-NEXT: v_fma_f32 v139, s4, v80, -v128
+ ; GCN-NEXT: v_pack_b32_f16 v143, v130, v131
; GCN-NEXT: v_cvt_f16_f32_e32 v138, v108
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[142:143], v[0:15]
; GCN-NEXT: v_exp_f32_e32 v80, v129
- ; GCN-NEXT: ds_read_b128 v[130:133], v197
- ; GCN-NEXT: s_waitcnt lgkmcnt(0)
- ; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: ds_read_b128 v[146:149], v197 offset:576
- ; GCN-NEXT: s_waitcnt lgkmcnt(0)
- ; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v139
; GCN-NEXT: v_cvt_f16_f32_e32 v139, v109
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[144:145], v[142:143], v[32:47]
- ; GCN-NEXT: v_fma_f32 v144, s4, v81, -v128
- ; GCN-NEXT: v_exp_f32_e32 v81, v125
- ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v144
- ; GCN-NEXT: v_pack_b32_f16 v144, v138, v139
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[136:137], v[142:143], v[16:31]
; GCN-NEXT: v_cvt_f16_f32_e32 v136, v110
; GCN-NEXT: v_fma_f32 v137, s4, v82, -v128
@@ -784,13 +775,25 @@
; GCN-NEXT: v_cvt_f16_f32_e32 v134, v111
; GCN-NEXT: v_mul_f32_e32 v156, 0x3fb8aa3b, v137
; GCN-NEXT: v_fma_f32 v137, s4, v83, -v128
- ; GCN-NEXT: v_mul_f32_e32 v157, 0x3fb8aa3b, v137
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[140:141], v[142:143], v[48:63]
; GCN-NEXT: v_exp_f32_e32 v83, v135
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[142:143], v[0:15]
+ ; GCN-NEXT: ds_read_b128 v[130:133], v197
+ ; GCN-NEXT: s_waitcnt lgkmcnt(0)
+ ; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: ds_read_b128 v[146:149], v197 offset:576
+ ; GCN-NEXT: s_waitcnt lgkmcnt(0)
+ ; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: v_mul_f32_e32 v157, 0x3fb8aa3b, v137
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[144:145], v[142:143], v[32:47]
; GCN-NEXT: v_pack_b32_f16 v145, v136, v134
; GCN-NEXT: ds_read_b128 v[134:137], v197 offset:1152
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: v_fma_f32 v144, s4, v81, -v128
+ ; GCN-NEXT: v_exp_f32_e32 v81, v125
+ ; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v144
+ ; GCN-NEXT: v_pack_b32_f16 v144, v138, v139
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[140:141], v[142:143], v[48:63]
; GCN-NEXT: ds_read_b128 v[138:141], v197 offset:1728
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
@@ -803,7 +806,6 @@
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GCN-NEXT: ds_write_b64 v200, v[150:151]
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[130:131], v[144:145], v[0:15]
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GCN-NEXT: ds_write_b64 v201, v[152:153]
@@ -811,19 +813,17 @@
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GCN-NEXT: ds_write_b64 v202, v[154:155]
; GCN-NEXT: v_fma_f32 v127, s4, v84, -v128
- ; GCN-NEXT: v_exp_f32_e32 v84, v129
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[130:131], v[144:145], v[0:15]
; GCN-NEXT: v_fma_f32 v130, s4, v85, -v128
- ; GCN-NEXT: v_cvt_f16_f32_e32 v126, v80
- ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v127
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[146:147], v[144:145], v[32:47]
; GCN-NEXT: v_exp_f32_e32 v85, v125
; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v130
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_load_dwordx2 v[130:131], v206, s[0:3], 0 offen sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_cvt_f16_f32_e32 v127, v81
- ; GCN-NEXT: v_pack_b32_f16 v126, v126, v127
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v126, v80
+ ; GCN-NEXT: v_exp_f32_e32 v84, v129
+ ; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v127
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[134:135], v[144:145], v[16:31]
; GCN-NEXT: v_fma_f32 v134, s4, v86, -v128
; GCN-NEXT: v_mul_f32_e32 v158, 0x3fb8aa3b, v134
@@ -833,27 +833,28 @@
; GCN-NEXT: buffer_load_dwordx2 v[142:143], v204, s[0:3], 0 offen sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v127, v81
+ ; GCN-NEXT: v_exp_f32_e32 v86, v156
+ ; GCN-NEXT: v_pack_b32_f16 v126, v126, v127
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[146:147], v[144:145], v[32:47]
; GCN-NEXT: buffer_load_dwordx2 v[146:147], v205, s[0:3], 0 offen sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: v_cvt_f16_f32_e32 v127, v82
- ; GCN-NEXT: v_exp_f32_e32 v86, v156
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[138:139], v[144:145], v[48:63]
- ; GCN-NEXT: v_cvt_f16_f32_e32 v138, v83
; GCN-NEXT: ;;#ASMSTART
; GCN-NEXT: s_waitcnt vmcnt(8)
; GCN-NEXT: ;;#ASMEND
- ; GCN-NEXT: v_fma_f32 v139, s4, v87, -v128
- ; GCN-NEXT: v_exp_f32_e32 v87, v157
- ; GCN-NEXT: v_pack_b32_f16 v127, v127, v138
- ; GCN-NEXT: v_fma_f32 v138, s4, v89, -v128
- ; GCN-NEXT: v_mul_f32_e32 v139, 0x3fb8aa3b, v139
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[126:127], v[0:15]
; GCN-NEXT: ; implicit-def: $sgpr0
; GCN-NEXT: v_perm_b32 v154, v135, v131, s5
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[138:139], v[144:145], v[48:63]
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v138, v83
; GCN-NEXT: v_perm_b32 v156, v135, v131, s7
; GCN-NEXT: v_fma_f32 v135, s4, v88, -v128
+ ; GCN-NEXT: v_fma_f32 v139, s4, v87, -v128
+ ; GCN-NEXT: v_pack_b32_f16 v127, v127, v138
+ ; GCN-NEXT: v_exp_f32_e32 v87, v157
; GCN-NEXT: v_perm_b32 v150, v134, v130, s5
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[126:127], v[0:15]
; GCN-NEXT: v_perm_b32 v152, v134, v130, s7
; GCN-NEXT: ds_read_b128 v[130:133], v198
; GCN-NEXT: s_waitcnt lgkmcnt(0)
@@ -870,13 +871,15 @@
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[148:149], v[126:127], v[32:47]
- ; GCN-NEXT: v_exp_f32_e32 v89, v125
; GCN-NEXT: v_pack_b32_f16 v146, v134, v135
; GCN-NEXT: v_cvt_f16_f32_e32 v134, v86
+ ; GCN-NEXT: v_fma_f32 v138, s4, v89, -v128
+ ; GCN-NEXT: v_exp_f32_e32 v89, v125
+ ; GCN-NEXT: v_mul_f32_e32 v139, 0x3fb8aa3b, v139
; GCN-NEXT: v_fma_f32 v135, s4, v90, -v128
; GCN-NEXT: v_mul_f32_e32 v125, 0x3fb8aa3b, v138
- ; GCN-NEXT: v_mul_f32_e32 v148, 0x3fb8aa3b, v135
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[136:137], v[126:127], v[16:31]
+ ; GCN-NEXT: v_mul_f32_e32 v148, 0x3fb8aa3b, v135
; GCN-NEXT: v_exp_f32_e32 v90, v158
; GCN-NEXT: v_mul_f32_e32 v158, 0x3fb8aa3b, v64
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[140:141], v[126:127], v[48:63]
@@ -898,46 +901,45 @@
; GCN-NEXT: v_mul_f32_e32 v129, 0x3fb8aa3b, v130
; GCN-NEXT: v_cvt_f16_f32_e32 v130, v89
; GCN-NEXT: v_fma_f32 v131, s4, v93, -v128
- ; GCN-NEXT: v_pack_b32_f16 v130, v126, v130
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[142:143], v[146:147], v[32:47]
; GCN-NEXT: v_exp_f32_e32 v93, v125
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[142:143], v[146:147], v[32:47]
+ ; GCN-NEXT: v_pack_b32_f16 v130, v126, v130
; GCN-NEXT: v_fma_f32 v126, s4, v94, -v128
; GCN-NEXT: v_cvt_f16_f32_e32 v125, v90
; GCN-NEXT: v_mul_f32_e32 v143, 0x3fb8aa3b, v126
; GCN-NEXT: v_cvt_f16_f32_e32 v126, v91
; GCN-NEXT: v_mul_f32_e32 v142, 0x3fb8aa3b, v131
- ; GCN-NEXT: v_fma_f32 v131, s4, v95, -v128
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[134:135], v[146:147], v[16:31]
; GCN-NEXT: v_exp_f32_e32 v94, v148
- ; GCN-NEXT: v_cvt_f16_f32_e32 v64, v93
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[138:139], v[146:147], v[48:63]
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[134:135], v[146:147], v[16:31]
+ ; GCN-NEXT: v_fma_f32 v131, s4, v95, -v128
; GCN-NEXT: v_exp_f32_e32 v95, v127
; GCN-NEXT: v_cvt_f16_f32_e32 v127, v92
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v64, v93
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[138:139], v[146:147], v[48:63]
; GCN-NEXT: v_mul_f32_e32 v138, 0x3fb8aa3b, v131
; GCN-NEXT: v_pack_b32_f16 v131, v125, v126
- ; GCN-NEXT: s_nop 1
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[130:131], v[0:15]
+ ; GCN-NEXT: v_exp_f32_e32 v126, v142
+ ; GCN-NEXT: v_pack_b32_f16 v142, v127, v64
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v64, v94
; GCN-NEXT: v_exp_f32_e32 v125, v129
- ; GCN-NEXT: ds_read_b128 v[132:135], v197
- ; GCN-NEXT: s_waitcnt lgkmcnt(0)
- ; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: ds_read_b128 v[146:149], v197 offset:576
- ; GCN-NEXT: s_waitcnt lgkmcnt(0)
- ; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: v_exp_f32_e32 v127, v143
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[144:145], v[130:131], v[32:47]
; GCN-NEXT: v_mul_f32_e32 v144, 0x3fb8aa3b, v65
; GCN-NEXT: v_fma_f32 v65, s4, v66, -v128
- ; GCN-NEXT: v_exp_f32_e32 v126, v142
- ; GCN-NEXT: v_pack_b32_f16 v142, v127, v64
- ; GCN-NEXT: v_cvt_f16_f32_e32 v64, v94
; GCN-NEXT: v_mul_f32_e32 v145, 0x3fb8aa3b, v65
; GCN-NEXT: v_cvt_f16_f32_e32 v65, v95
; GCN-NEXT: v_fma_f32 v66, s4, v67, -v128
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[136:137], v[130:131], v[16:31]
- ; GCN-NEXT: v_exp_f32_e32 v127, v143
+ ; GCN-NEXT: v_exp_f32_e32 v129, v138
; GCN-NEXT: v_pack_b32_f16 v143, v64, v65
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[130:131], v[0:15]
+ ; GCN-NEXT: ds_read_b128 v[132:135], v197
+ ; GCN-NEXT: s_waitcnt lgkmcnt(0)
+ ; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: ds_read_b128 v[146:149], v197 offset:576
+ ; GCN-NEXT: s_waitcnt lgkmcnt(0)
+ ; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[136:137], v[130:131], v[16:31]
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[140:141], v[130:131], v[48:63]
- ; GCN-NEXT: v_exp_f32_e32 v129, v138
; GCN-NEXT: v_mul_f32_e32 v141, 0x3fb8aa3b, v66
; GCN-NEXT: ds_read_b128 v[64:67], v197 offset:1152
; GCN-NEXT: s_waitcnt lgkmcnt(0)
@@ -945,127 +947,117 @@
; GCN-NEXT: ds_read_b128 v[136:139], v197 offset:1728
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: v_exp_f32_e32 v130, v158
+ ; GCN-NEXT: v_exp_f32_e32 v131, v144
; GCN-NEXT: ;;#ASMSTART
; GCN-NEXT: s_waitcnt vmcnt(8)
; GCN-NEXT: ;;#ASMEND
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: ds_write_b64 v199, v[150:151]
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[142:143], v[0:15]
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v132, v125
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GCN-NEXT: ds_write_b64 v200, v[152:153]
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[132:133], v[142:143], v[0:15]
- ; GCN-NEXT: v_cvt_f16_f32_e32 v132, v125
- ; GCN-NEXT: v_exp_f32_e32 v130, v158
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GCN-NEXT: ds_write_b64 v201, v[154:155]
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GCN-NEXT: ds_write_b64 v202, v[156:157]
+ ; GCN-NEXT: v_exp_f32_e32 v133, v141
; GCN-NEXT: ;;#ASMSTART
; GCN-NEXT: s_waitcnt vmcnt(8)
; GCN-NEXT: ;;#ASMEND
+ ; GCN-NEXT: v_mul_f32_e32 v144, 0x3fb8aa3b, v69
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[146:147], v[142:143], v[32:47]
; GCN-NEXT: v_mul_f32_e32 v146, 0x3fb8aa3b, v68
; GCN-NEXT: v_cvt_f16_f32_e32 v68, v126
- ; GCN-NEXT: v_exp_f32_e32 v131, v144
- ; GCN-NEXT: v_mul_f32_e32 v144, 0x3fb8aa3b, v69
; GCN-NEXT: v_fma_f32 v69, s4, v71, -v128
; GCN-NEXT: v_pack_b32_f16 v140, v132, v68
; GCN-NEXT: v_cvt_f16_f32_e32 v68, v129
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[64:65], v[142:143], v[16:31]
- ; GCN-NEXT: v_cvt_f16_f32_e32 v64, v127
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[136:137], v[142:143], v[48:63]
; GCN-NEXT: v_exp_f32_e32 v132, v145
- ; GCN-NEXT: v_fma_f32 v65, s4, v70, -v128
- ; GCN-NEXT: v_mul_f32_e32 v65, 0x3fb8aa3b, v65
; GCN-NEXT: v_fma_f32 v145, s4, v73, -v128
+ ; GCN-NEXT: v_exp_f32_e32 v73, v144
; GCN-NEXT: v_mul_f32_e32 v147, 0x3fb8aa3b, v145
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[136:137], v[142:143], v[48:63]
- ; GCN-NEXT: v_exp_f32_e32 v133, v141
- ; GCN-NEXT: v_mul_f32_e32 v142, 0x3fb8aa3b, v69
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[64:65], v[142:143], v[16:31]
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v64, v127
+ ; GCN-NEXT: v_fma_f32 v143, s4, v72, -v128
+ ; GCN-NEXT: v_exp_f32_e32 v72, v146
+ ; GCN-NEXT: v_mul_f32_e32 v146, 0x3fb8aa3b, v143
; GCN-NEXT: v_pack_b32_f16 v141, v64, v68
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v64, v130
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v143, v131
+ ; GCN-NEXT: v_fma_f32 v65, s4, v70, -v128
+ ; GCN-NEXT: v_mul_f32_e32 v142, 0x3fb8aa3b, v69
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: ds_read_b128 v[68:71], v198
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_fma_f32 v143, s4, v72, -v128
- ; GCN-NEXT: v_cvt_f16_f32_e32 v64, v130
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[134:135], v[140:141], v[0:15]
- ; GCN-NEXT: v_exp_f32_e32 v72, v146
- ; GCN-NEXT: v_mul_f32_e32 v146, 0x3fb8aa3b, v143
- ; GCN-NEXT: v_cvt_f16_f32_e32 v143, v131
; GCN-NEXT: ds_read_b128 v[134:137], v198 offset:576
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: v_pack_b32_f16 v64, v64, v143
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[148:149], v[140:141], v[32:47]
- ; GCN-NEXT: v_exp_f32_e32 v73, v144
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[66:67], v[140:141], v[16:31]
- ; GCN-NEXT: v_cvt_f16_f32_e32 v66, v132
- ; GCN-NEXT: v_fma_f32 v67, s4, v74, -v128
- ; GCN-NEXT: v_exp_f32_e32 v74, v65
- ; GCN-NEXT: v_cvt_f16_f32_e32 v65, v133
- ; GCN-NEXT: v_mul_f32_e32 v67, 0x3fb8aa3b, v67
- ; GCN-NEXT: v_pack_b32_f16 v65, v66, v65
+ ; GCN-NEXT: v_mul_f32_e32 v65, 0x3fb8aa3b, v65
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[138:139], v[140:141], v[48:63]
; GCN-NEXT: v_fma_f32 v138, s4, v75, -v128
; GCN-NEXT: v_exp_f32_e32 v75, v142
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[148:149], v[140:141], v[32:47]
; GCN-NEXT: v_mul_f32_e32 v148, 0x3fb8aa3b, v138
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[66:67], v[140:141], v[16:31]
; GCN-NEXT: ds_read_b128 v[138:141], v198 offset:1152
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: ds_read_b128 v[142:145], v198 offset:1728
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v66, v132
+ ; GCN-NEXT: v_fma_f32 v67, s4, v74, -v128
+ ; GCN-NEXT: v_exp_f32_e32 v74, v65
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v65, v133
+ ; GCN-NEXT: v_mul_f32_e32 v67, 0x3fb8aa3b, v67
+ ; GCN-NEXT: v_pack_b32_f16 v65, v66, v65
; GCN-NEXT: v_cvt_f16_f32_e32 v66, v72
+ ; GCN-NEXT: s_nop 0
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[68:69], v[64:65], v[0:15]
; GCN-NEXT: v_fma_f32 v68, s4, v76, -v128
; GCN-NEXT: v_exp_f32_e32 v76, v146
; GCN-NEXT: v_mul_f32_e32 v146, 0x3fb8aa3b, v68
; GCN-NEXT: v_cvt_f16_f32_e32 v68, v73
; GCN-NEXT: v_fma_f32 v69, s4, v77, -v128
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[134:135], v[64:65], v[32:47]
; GCN-NEXT: v_exp_f32_e32 v77, v147
+ ; GCN-NEXT: v_mul_f32_e32 v147, 0x3fb8aa3b, v69
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[142:143], v[64:65], v[48:63]
+ ; GCN-NEXT: v_exp_f32_e32 v142, v146
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[134:135], v[64:65], v[32:47]
; GCN-NEXT: v_pack_b32_f16 v134, v66, v68
- ; GCN-NEXT: v_fma_f32 v68, s4, v78, -v128
; GCN-NEXT: v_cvt_f16_f32_e32 v66, v74
- ; GCN-NEXT: v_mul_f32_e32 v147, 0x3fb8aa3b, v69
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[138:139], v[64:65], v[16:31]
+ ; GCN-NEXT: v_fma_f32 v68, s4, v78, -v128
; GCN-NEXT: v_exp_f32_e32 v78, v67
- ; GCN-NEXT: v_mul_f32_e32 v138, 0x3fb8aa3b, v68
- ; GCN-NEXT: v_cvt_f16_f32_e32 v139, v76
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[142:143], v[64:65], v[48:63]
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[138:139], v[64:65], v[16:31]
; GCN-NEXT: v_cvt_f16_f32_e32 v64, v75
; GCN-NEXT: v_fma_f32 v65, s4, v79, -v128
- ; GCN-NEXT: v_exp_f32_e32 v79, v148
+ ; GCN-NEXT: v_mul_f32_e32 v138, 0x3fb8aa3b, v68
; GCN-NEXT: v_mul_f32_e32 v128, 0x3fb8aa3b, v65
; GCN-NEXT: v_pack_b32_f16 v135, v66, v64
- ; GCN-NEXT: s_nop 1
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[70:71], v[134:135], v[0:15]
- ; GCN-NEXT: v_exp_f32_e32 v142, v146
- ; GCN-NEXT: ds_read_b128 v[68:71], v197
- ; GCN-NEXT: s_waitcnt lgkmcnt(0)
- ; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: ds_read_b128 v[64:67], v197 offset:576
- ; GCN-NEXT: s_waitcnt lgkmcnt(0)
- ; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[136:137], v[134:135], v[32:47]
- ; GCN-NEXT: v_exp_f32_e32 v137, v147
- ; GCN-NEXT: v_cvt_f16_f32_e32 v136, v77
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[140:141], v[134:135], v[16:31]
; GCN-NEXT: v_exp_f32_e32 v138, v138
- ; GCN-NEXT: v_cvt_f16_f32_e32 v140, v78
+ ; GCN-NEXT: v_exp_f32_e32 v79, v148
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[144:145], v[134:135], v[48:63]
- ; GCN-NEXT: s_nop 10
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v139, v76
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[136:137], v[134:135], v[32:47]
+ ; GCN-NEXT: v_exp_f32_e32 v137, v147
+ ; GCN-NEXT: s_nop 7
; GCN-NEXT: v_exp_f32_e32 v52, v128
- ; GCN-NEXT: v_cvt_f16_f32_e32 v50, v137
; GCN-NEXT: v_cvt_f16_f32_e32 v51, v142
; GCN-NEXT: v_cvt_f16_f32_e32 v54, v138
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v50, v137
; GCN-NEXT: v_cvt_f16_f32_e32 v53, v52
- ; GCN-NEXT: v_cvt_f16_f32_e32 v49, v79
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v136, v77
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[70:71], v[134:135], v[0:15]
; GCN-NEXT: v_pack_b32_f16 v50, v51, v50
- ; GCN-NEXT: v_pack_b32_f16 v48, v139, v136
; GCN-NEXT: v_pack_b32_f16 v51, v54, v53
; GCN-NEXT: v_add_f32_e32 v53, 0, v113
; GCN-NEXT: v_add_f32_e32 v53, v114, v53
@@ -1097,14 +1089,21 @@
; GCN-NEXT: v_add_f32_e32 v53, v111, v53
; GCN-NEXT: v_add_f32_e32 v53, v80, v53
; GCN-NEXT: v_add_f32_e32 v53, v81, v53
+ ; GCN-NEXT: ds_read_b128 v[68:71], v197
+ ; GCN-NEXT: s_waitcnt lgkmcnt(0)
+ ; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: v_add_f32_e32 v53, v82, v53
; GCN-NEXT: v_add_f32_e32 v53, v83, v53
; GCN-NEXT: v_add_f32_e32 v53, v84, v53
; GCN-NEXT: v_add_f32_e32 v53, v85, v53
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[16:31], v[140:141], v[134:135], v[16:31]
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v140, v78
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v49, v79
; GCN-NEXT: v_add_f32_e32 v53, v86, v53
; GCN-NEXT: v_add_f32_e32 v53, v87, v53
; GCN-NEXT: v_add_f32_e32 v53, v88, v53
; GCN-NEXT: v_add_f32_e32 v53, v89, v53
+ ; GCN-NEXT: v_pack_b32_f16 v48, v139, v136
; GCN-NEXT: v_pack_b32_f16 v49, v140, v49
; GCN-NEXT: v_add_f32_e32 v53, v90, v53
; GCN-NEXT: v_add_f32_e32 v53, v91, v53
@@ -1118,8 +1117,11 @@
; GCN-NEXT: v_add_f32_e32 v53, v127, v53
; GCN-NEXT: v_add_f32_e32 v53, v129, v53
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[0:15], v[70:71], v[50:51], v[0:15]
+ ; GCN-NEXT: ds_read_b128 v[64:67], v197 offset:576
+ ; GCN-NEXT: s_waitcnt lgkmcnt(0)
+ ; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[64:65], v[48:49], v[32:47]
- ; GCN-NEXT: s_nop 9
+ ; GCN-NEXT: s_nop 6
; GCN-NEXT: v_add_f32_e32 v0, v130, v53
; GCN-NEXT: v_add_f32_e32 v0, v131, v0
; GCN-NEXT: v_add_f32_e32 v0, v132, v0
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.small.mir b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.small.mir
index 0a8d7acd187fc..713e8f4c791dd 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.small.mir
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.iglp.opt.exp.small.mir
@@ -221,51 +221,47 @@
; GCN-NEXT: v_max_f32_e32 v70, v70, v70
; GCN-NEXT: v_max_f32_e32 v72, v81, v70
; GCN-NEXT: v_fma_f32 v16, s4, v16, -v72
+ ; GCN-NEXT: v_fma_f32 v17, s4, v17, -v72
; GCN-NEXT: v_fma_f32 v18, s4, v18, -v72
; GCN-NEXT: v_fma_f32 v19, s4, v19, -v72
; GCN-NEXT: v_mul_f32_e32 v16, 0x3fb8aa3b, v16
+ ; GCN-NEXT: v_mul_f32_e32 v17, 0x3fb8aa3b, v17
; GCN-NEXT: v_mul_f32_e32 v18, 0x3fb8aa3b, v18
; GCN-NEXT: v_mul_f32_e32 v19, 0x3fb8aa3b, v19
- ; GCN-NEXT: v_fma_f32 v17, s4, v17, -v72
- ; GCN-NEXT: v_fma_f32 v20, s4, v20, -v72
- ; GCN-NEXT: v_fma_f32 v21, s4, v21, -v72
- ; GCN-NEXT: v_fma_f32 v22, s4, v22, -v72
- ; GCN-NEXT: v_fma_f32 v23, s4, v23, -v72
; GCN-NEXT: v_exp_f32_e32 v73, v16
+ ; GCN-NEXT: v_exp_f32_e32 v17, v17
; GCN-NEXT: v_exp_f32_e32 v74, v18
; GCN-NEXT: v_exp_f32_e32 v75, v19
- ; GCN-NEXT: v_mul_f32_e32 v20, 0x3fb8aa3b, v20
+ ; GCN-NEXT: v_fma_f32 v21, s4, v21, -v72
+ ; GCN-NEXT: v_fma_f32 v22, s4, v22, -v72
; GCN-NEXT: v_mul_f32_e32 v21, 0x3fb8aa3b, v21
; GCN-NEXT: v_mul_f32_e32 v22, 0x3fb8aa3b, v22
- ; GCN-NEXT: v_exp_f32_e32 v80, v20
; GCN-NEXT: v_cvt_f16_f32_e32 v16, v73
- ; GCN-NEXT: v_fma_f32 v18, s4, v24, -v72
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v19, v17
; GCN-NEXT: v_exp_f32_e32 v81, v21
; GCN-NEXT: v_cvt_f16_f32_e32 v21, v74
- ; GCN-NEXT: v_fma_f32 v20, s4, v25, -v72
; GCN-NEXT: v_exp_f32_e32 v82, v22
; GCN-NEXT: v_cvt_f16_f32_e32 v22, v75
- ; GCN-NEXT: v_mul_f32_e32 v17, 0x3fb8aa3b, v17
+ ; GCN-NEXT: v_fma_f32 v20, s4, v20, -v72
+ ; GCN-NEXT: v_fma_f32 v23, s4, v23, -v72
+ ; GCN-NEXT: v_mul_f32_e32 v20, 0x3fb8aa3b, v20
+ ; GCN-NEXT: v_fma_f32 v18, s4, v24, -v72
+ ; GCN-NEXT: v_exp_f32_e32 v80, v20
+ ; GCN-NEXT: v_fma_f32 v20, s4, v25, -v72
+ ; GCN-NEXT: v_sub_f32_e32 v24, v67, v72
; GCN-NEXT: v_mul_f32_e32 v23, 0x3fb8aa3b, v23
- ; GCN-NEXT: v_fma_f32 v26, s4, v26, -v72
+ ; GCN-NEXT: v_mul_f32_e32 v24, 0x3fb8aa3b, v24
; GCN-NEXT: v_pack_b32_f16 v71, v21, v22
- ; GCN-NEXT: v_mul_f32_e32 v22, 0x3fb8aa3b, v18
- ; GCN-NEXT: v_sub_f32_e32 v24, v67, v72
+ ; GCN-NEXT: v_pack_b32_f16 v70, v16, v19
; GCN-NEXT: v_exp_f32_e32 v83, v23
- ; GCN-NEXT: v_fma_f32 v67, s4, v27, -v72
- ; GCN-NEXT: v_exp_f32_e32 v85, v22
- ; GCN-NEXT: v_exp_f32_e32 v17, v17
- ; GCN-NEXT: v_mul_f32_e32 v24, 0x3fb8aa3b, v24
+ ; GCN-NEXT: v_mul_f32_e32 v22, 0x3fb8aa3b, v18
; GCN-NEXT: v_mul_f32_e32 v23, 0x3fb8aa3b, v20
- ; GCN-NEXT: v_cvt_f16_f32_e32 v19, v17
- ; GCN-NEXT: v_fma_f32 v87, s4, v29, -v72
- ; GCN-NEXT: v_exp_f32_e32 v88, v23
- ; GCN-NEXT: v_fma_f32 v0, s4, v0, -v72
- ; GCN-NEXT: v_pack_b32_f16 v70, v16, v19
; GCN-NEXT: ds_read_b128 v[18:21], v84
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
; GCN-NEXT: v_exp_f32_e32 v16, v24
+ ; GCN-NEXT: v_exp_f32_e32 v85, v22
+ ; GCN-NEXT: v_exp_f32_e32 v88, v23
; GCN-NEXT: ds_read_b128 v[22:25], v84 offset:576
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
@@ -286,22 +282,24 @@
; GCN-NEXT: v_pk_mul_f32 v[44:45], v[44:45], v[16:17] op_sel_hi:[1,0]
; GCN-NEXT: v_pk_mul_f32 v[46:47], v[46:47], v[16:17] op_sel_hi:[1,0]
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[18:19], v[70:71], v[48:63]
+ ; GCN-NEXT: v_fma_f32 v26, s4, v26, -v72
; GCN-NEXT: v_add_f32_e32 v18, 0, v73
- ; GCN-NEXT: v_cvt_f16_f32_e32 v89, v83
- ; GCN-NEXT: v_fma_f32 v73, s4, v28, -v72
; GCN-NEXT: v_cvt_f16_f32_e32 v19, v80
- ; GCN-NEXT: v_fma_f32 v1, s4, v1, -v72
- ; GCN-NEXT: v_perm_b32 v90, v69, v65, s2
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[22:23], v[70:71], v[32:47]
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v86, v81
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v89, v83
; GCN-NEXT: v_add_f32_e32 v17, v17, v18
; GCN-NEXT: v_mul_f32_e32 v18, 0x3fb8aa3b, v26
- ; GCN-NEXT: v_cvt_f16_f32_e32 v86, v81
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[22:23], v[70:71], v[32:47]
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v22, v82
+ ; GCN-NEXT: v_fma_f32 v67, s4, v27, -v72
; GCN-NEXT: v_fma_f32 v23, s4, v30, -v72
; GCN-NEXT: v_exp_f32_e32 v30, v18
- ; GCN-NEXT: v_cvt_f16_f32_e32 v22, v82
; GCN-NEXT: v_fma_f32 v18, s4, v31, -v72
; GCN-NEXT: v_perm_b32 v31, v68, v64, s2
+ ; GCN-NEXT: v_fma_f32 v73, s4, v28, -v72
+ ; GCN-NEXT: v_fma_f32 v87, s4, v29, -v72
; GCN-NEXT: v_perm_b32 v64, v68, v64, s3
+ ; GCN-NEXT: v_perm_b32 v90, v69, v65, s2
; GCN-NEXT: v_perm_b32 v65, v69, v65, s3
; GCN-NEXT: ds_read_b128 v[26:29], v91
; GCN-NEXT: s_waitcnt lgkmcnt(0)
@@ -326,50 +324,46 @@
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GCN-NEXT: ds_write_b32 v78, v90
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[20:21], v[18:19], v[48:63]
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v20, v85
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v21, v88
; GCN-NEXT: buffer_wbl2 sc0 sc1
; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
; GCN-NEXT: ds_write_b32 v79, v65
; GCN-NEXT: v_mul_f32_e32 v64, 0x3fb8aa3b, v73
; GCN-NEXT: v_mul_f32_e32 v65, 0x3fb8aa3b, v87
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[20:21], v[18:19], v[48:63]
- ; GCN-NEXT: v_add_f32_e32 v17, v74, v17
- ; GCN-NEXT: v_cvt_f16_f32_e32 v20, v85
- ; GCN-NEXT: v_fma_f32 v2, s4, v2, -v72
- ; GCN-NEXT: v_exp_f32_e32 v22, v64
- ; GCN-NEXT: v_cvt_f16_f32_e32 v21, v88
- ; GCN-NEXT: v_exp_f32_e32 v64, v65
; GCN-NEXT: v_mul_f32_e32 v23, 0x3fb8aa3b, v23
+ ; GCN-NEXT: v_fma_f32 v0, s4, v0, -v72
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[24:25], v[18:19], v[32:47]
- ; GCN-NEXT: v_add_f32_e32 v17, v75, v17
; GCN-NEXT: v_cvt_f16_f32_e32 v18, v30
- ; GCN-NEXT: v_fma_f32 v24, s4, v3, -v72
- ; GCN-NEXT: v_exp_f32_e32 v23, v23
; GCN-NEXT: v_cvt_f16_f32_e32 v19, v31
+ ; GCN-NEXT: v_fma_f32 v1, s4, v1, -v72
+ ; GCN-NEXT: v_exp_f32_e32 v22, v64
+ ; GCN-NEXT: v_exp_f32_e32 v64, v65
+ ; GCN-NEXT: v_exp_f32_e32 v23, v23
+ ; GCN-NEXT: v_fma_f32 v24, s4, v3, -v72
+ ; GCN-NEXT: v_exp_f32_e32 v25, v67
; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v0
; GCN-NEXT: v_mul_f32_e32 v65, 0x3fb8aa3b, v1
; GCN-NEXT: v_pack_b32_f16 v0, v20, v21
; GCN-NEXT: v_pack_b32_f16 v1, v18, v19
- ; GCN-NEXT: v_fma_f32 v6, s4, v6, -v72
- ; GCN-NEXT: v_exp_f32_e32 v25, v67
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[26:27], v[0:1], v[48:63]
- ; GCN-NEXT: v_add_f32_e32 v17, v80, v17
; GCN-NEXT: v_cvt_f16_f32_e32 v18, v22
+ ; GCN-NEXT: v_fma_f32 v67, s4, v5, -v72
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[26:27], v[0:1], v[48:63]
; GCN-NEXT: v_fma_f32 v26, s4, v4, -v72
- ; GCN-NEXT: v_exp_f32_e32 v27, v3
; GCN-NEXT: v_cvt_f16_f32_e32 v4, v64
- ; GCN-NEXT: v_fma_f32 v67, s4, v5, -v72
- ; GCN-NEXT: v_exp_f32_e32 v65, v65
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[68:69], v[0:1], v[32:47]
- ; GCN-NEXT: v_mul_f32_e32 v2, 0x3fb8aa3b, v2
- ; GCN-NEXT: v_add_f32_e32 v17, v81, v17
; GCN-NEXT: v_cvt_f16_f32_e32 v5, v23
- ; GCN-NEXT: v_fma_f32 v7, s4, v7, -v72
- ; GCN-NEXT: v_exp_f32_e32 v68, v2
; GCN-NEXT: v_cvt_f16_f32_e32 v19, v25
+ ; GCN-NEXT: v_fma_f32 v2, s4, v2, -v72
+ ; GCN-NEXT: v_mul_f32_e32 v2, 0x3fb8aa3b, v2
; GCN-NEXT: ;;#ASMSTART
; GCN-NEXT: s_waitcnt vmcnt(8)
; GCN-NEXT: ;;#ASMEND
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[68:69], v[0:1], v[32:47]
; GCN-NEXT: v_mul_f32_e32 v24, 0x3fb8aa3b, v24
+ ; GCN-NEXT: v_exp_f32_e32 v27, v3
+ ; GCN-NEXT: v_exp_f32_e32 v65, v65
+ ; GCN-NEXT: v_exp_f32_e32 v68, v2
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: ds_read_b128 v[0:3], v84
; GCN-NEXT: s_waitcnt lgkmcnt(0)
@@ -380,42 +374,59 @@
; GCN-NEXT: ds_read_b128 v[18:21], v84 offset:576
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
- ; GCN-NEXT: v_mul_f32_e32 v26, 0x3fb8aa3b, v26
- ; GCN-NEXT: v_mul_f32_e32 v67, 0x3fb8aa3b, v67
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[28:29], v[4:5], v[48:63]
- ; GCN-NEXT: v_add_f32_e32 v17, v82, v17
+ ; GCN-NEXT: v_add_f32_e32 v17, v74, v17
; GCN-NEXT: v_cvt_f16_f32_e32 v28, v27
- ; GCN-NEXT: v_exp_f32_e32 v26, v26
; GCN-NEXT: v_cvt_f16_f32_e32 v29, v65
- ; GCN-NEXT: v_fma_f32 v10, s4, v10, -v72
- ; GCN-NEXT: v_exp_f32_e32 v67, v67
- ; GCN-NEXT: v_mul_f32_e32 v6, 0x3fb8aa3b, v6
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v69, v24
+ ; GCN-NEXT: v_add_f32_e32 v17, v75, v17
+ ; GCN-NEXT: v_add_f32_e32 v17, v80, v17
+ ; GCN-NEXT: v_fma_f32 v6, s4, v6, -v72
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[70:71], v[4:5], v[32:47]
- ; GCN-NEXT: v_add_f32_e32 v17, v83, v17
; GCN-NEXT: v_cvt_f16_f32_e32 v5, v68
- ; GCN-NEXT: v_exp_f32_e32 v6, v6
- ; GCN-NEXT: v_cvt_f16_f32_e32 v69, v24
+ ; GCN-NEXT: v_fma_f32 v7, s4, v7, -v72
+ ; GCN-NEXT: v_add_f32_e32 v17, v81, v17
+ ; GCN-NEXT: v_mul_f32_e32 v26, 0x3fb8aa3b, v26
+ ; GCN-NEXT: v_mul_f32_e32 v67, 0x3fb8aa3b, v67
+ ; GCN-NEXT: v_mul_f32_e32 v6, 0x3fb8aa3b, v6
; GCN-NEXT: v_mul_f32_e32 v7, 0x3fb8aa3b, v7
+ ; GCN-NEXT: v_add_f32_e32 v17, v82, v17
+ ; GCN-NEXT: v_exp_f32_e32 v26, v26
+ ; GCN-NEXT: v_exp_f32_e32 v67, v67
+ ; GCN-NEXT: v_exp_f32_e32 v6, v6
; GCN-NEXT: v_exp_f32_e32 v7, v7
; GCN-NEXT: v_pack_b32_f16 v4, v28, v29
; GCN-NEXT: v_pack_b32_f16 v5, v5, v69
- ; GCN-NEXT: ; implicit-def: $sgpr2
- ; GCN-NEXT: s_nop 1
+ ; GCN-NEXT: v_add_f32_e32 v17, v83, v17
+ ; GCN-NEXT: v_fma_f32 v10, s4, v10, -v72
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[0:1], v[4:5], v[48:63]
; GCN-NEXT: v_add_f32_e32 v0, v85, v17
; GCN-NEXT: v_cvt_f16_f32_e32 v17, v26
; GCN-NEXT: v_cvt_f16_f32_e32 v28, v67
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v1, v6
+ ; GCN-NEXT: v_fma_f32 v8, s4, v8, -v72
+ ; GCN-NEXT: v_mul_f32_e32 v8, 0x3fb8aa3b, v8
+ ; GCN-NEXT: v_exp_f32_e32 v8, v8
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[18:19], v[4:5], v[32:47]
; GCN-NEXT: v_add_f32_e32 v4, v88, v0
; GCN-NEXT: v_mul_f32_e32 v0, 0x3fb8aa3b, v10
- ; GCN-NEXT: v_cvt_f16_f32_e32 v1, v6
; GCN-NEXT: v_exp_f32_e32 v10, v0
; GCN-NEXT: v_cvt_f16_f32_e32 v0, v7
+ ; GCN-NEXT: v_fma_f32 v12, s4, v12, -v72
+ ; GCN-NEXT: v_fma_f32 v5, s4, v14, -v72
+ ; GCN-NEXT: v_fma_f32 v9, s4, v9, -v72
; GCN-NEXT: v_pack_b32_f16 v1, v1, v0
; GCN-NEXT: v_pack_b32_f16 v0, v17, v28
- ; GCN-NEXT: s_nop 1
+ ; GCN-NEXT: v_fma_f32 v11, s4, v11, -v72
+ ; GCN-NEXT: v_fma_f32 v13, s4, v13, -v72
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[2:3], v[0:1], v[48:63]
; GCN-NEXT: v_add_f32_e32 v2, v30, v4
+ ; GCN-NEXT: v_fma_f32 v14, s4, v15, -v72
+ ; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v12
+ ; GCN-NEXT: v_mul_f32_e32 v9, 0x3fb8aa3b, v9
+ ; GCN-NEXT: v_mul_f32_e32 v11, 0x3fb8aa3b, v11
+ ; GCN-NEXT: v_exp_f32_e32 v12, v3
+ ; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v13
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[20:21], v[0:1], v[32:47]
; GCN-NEXT: v_add_f32_e32 v0, v31, v2
; GCN-NEXT: v_add_f32_e32 v0, v22, v0
@@ -423,40 +434,26 @@
; GCN-NEXT: v_add_f32_e32 v0, v23, v0
; GCN-NEXT: v_add_f32_e32 v0, v25, v0
; GCN-NEXT: v_add_f32_e32 v0, v27, v0
- ; GCN-NEXT: v_fma_f32 v8, s4, v8, -v72
; GCN-NEXT: v_add_f32_e32 v0, v65, v0
- ; GCN-NEXT: v_fma_f32 v9, s4, v9, -v72
- ; GCN-NEXT: v_mul_f32_e32 v8, 0x3fb8aa3b, v8
; GCN-NEXT: v_add_f32_e32 v0, v68, v0
- ; GCN-NEXT: v_fma_f32 v11, s4, v11, -v72
- ; GCN-NEXT: v_mul_f32_e32 v9, 0x3fb8aa3b, v9
- ; GCN-NEXT: v_fma_f32 v12, s4, v12, -v72
- ; GCN-NEXT: v_fma_f32 v13, s4, v13, -v72
- ; GCN-NEXT: v_exp_f32_e32 v8, v8
; GCN-NEXT: v_add_f32_e32 v0, v24, v0
- ; GCN-NEXT: v_fma_f32 v5, s4, v14, -v72
- ; GCN-NEXT: v_exp_f32_e32 v9, v9
; GCN-NEXT: v_add_f32_e32 v0, v26, v0
; GCN-NEXT: v_add_f32_e32 v0, v67, v0
- ; GCN-NEXT: v_fma_f32 v14, s4, v15, -v72
- ; GCN-NEXT: v_mul_f32_e32 v11, 0x3fb8aa3b, v11
- ; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v12
; GCN-NEXT: v_mul_f32_e32 v1, 0x3fb8aa3b, v5
; GCN-NEXT: v_add_f32_e32 v0, v6, v0
- ; GCN-NEXT: v_exp_f32_e32 v11, v11
- ; GCN-NEXT: v_cvt_f16_f32_e32 v4, v8
- ; GCN-NEXT: v_exp_f32_e32 v12, v3
- ; GCN-NEXT: v_mul_f32_e32 v3, 0x3fb8aa3b, v13
; GCN-NEXT: v_exp_f32_e32 v17, v1
; GCN-NEXT: v_mul_f32_e32 v1, 0x3fb8aa3b, v14
; GCN-NEXT: v_add_f32_e32 v0, v7, v0
- ; GCN-NEXT: v_cvt_f16_f32_e32 v13, v9
+ ; GCN-NEXT: v_exp_f32_e32 v9, v9
+ ; GCN-NEXT: v_exp_f32_e32 v11, v11
; GCN-NEXT: v_exp_f32_e32 v15, v3
; GCN-NEXT: v_exp_f32_e32 v18, v1
; GCN-NEXT: v_add_f32_e32 v6, v8, v0
; GCN-NEXT: ds_read_b128 v[0:3], v91
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: buffer_inv sc0 sc1
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v4, v8
+ ; GCN-NEXT: v_cvt_f16_f32_e32 v13, v9
; GCN-NEXT: v_cvt_f16_f32_e32 v5, v10
; GCN-NEXT: v_cvt_f16_f32_e32 v14, v11
; GCN-NEXT: v_add_f32_e32 v6, v9, v6
@@ -482,13 +479,14 @@
; GCN-NEXT: s_waitcnt vmcnt(8)
; GCN-NEXT: ;;#ASMEND
; GCN-NEXT: v_mov_b32_e32 v4, 0
- ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[6:7], v[0:1], v[32:47]
+ ; GCN-NEXT: ; implicit-def: $sgpr2
; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[48:63], v[2:3], v[0:1], v[48:63]
; GCN-NEXT: v_add_f32_e32 v2, v18, v11
; GCN-NEXT: ds_bpermute_b32 v3, v66, v2
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: v_add_f32_e32 v2, v2, v3
; GCN-NEXT: ds_bpermute_b32 v3, v66, v2
+ ; GCN-NEXT: v_mfma_f32_32x32x8_f16 v[32:47], v[6:7], v[0:1], v[32:47]
; GCN-NEXT: s_waitcnt lgkmcnt(0)
; GCN-NEXT: v_cndmask_b32_e64 v2, v3, v2, s[0:1]
; GCN-NEXT: v_fmac_f32_e32 v2, v4, v16
More information about the llvm-commits
mailing list