[llvm] [AMDGPU] Post-RA AGPR copy elimination pass (PR #153901)
via llvm-commits
llvm-commits at lists.llvm.org
Fri Aug 15 16:18:46 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Lucas Ramirez (lucas-rami)
<details>
<summary>Changes</summary>
This introduces a post-RA pass to eliminate useless AGPR to AVGPR copies, effectively folding them into MFMAs' OpA and OpB when possible. These cross-class copies appear during register allocation when splitting live ranges lead to VGPR to AVGPR inflation on some parts of the original live range. RA is then unable to fold at least part of these copies.
The pass eliminates copies in the following kind of situations, making the MFMA use `a[0:3]` as `<opA>` instead.
```
v_accvgpr_read_b32 v0, a0
v_accvgpr_read_b32 v1, a1
v_accvgpr_read_b32 v2, a2
v_accvgpr_read_b32 v3, a3
v_mfma_f32_16x16x32_fp8_fp8 <dst>, v[0:3], <opB>, <opC>
```
This is still somewhat WIP and for now just seeking any feedback. Two questions I have in particular.
1. Should what this pass does be integrated in the `AMDGPURewriteAGPRCopyMFMA` pass instead of being its own thing?
2. This may be a good short-term solution to get rid of this kind of copies, but the pass is somewhat addressing the symptom rather than the cause of the issue, which is in the register allocator. Is this something we should strive to integrate at all?
---
Patch is 84.83 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/153901.diff
10 Files Affected:
- (modified) llvm/lib/Target/AMDGPU/AMDGPU.h (+11)
- (added) llvm/lib/Target/AMDGPU/AMDGPUEliminateAGPRToVGPRCopy.cpp (+246)
- (modified) llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def (+1)
- (modified) llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp (+2)
- (modified) llvm/lib/Target/AMDGPU/CMakeLists.txt (+1)
- (modified) llvm/test/CodeGen/AMDGPU/llc-pipeline.ll (+4)
- (modified) llvm/test/CodeGen/AMDGPU/mfma-loop.ll (+564-382)
- (modified) llvm/test/CodeGen/AMDGPU/sgpr-regalloc-flags.ll (+4)
- (modified) llvm/test/CodeGen/AMDGPU/spill-agpr.ll (+8-30)
- (modified) llvm/utils/gn/secondary/llvm/lib/Target/AMDGPU/BUILD.gn (+1)
``````````diff
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.h b/llvm/lib/Target/AMDGPU/AMDGPU.h
index 007b481f84960..b83e62f4ea57c 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.h
@@ -541,6 +541,17 @@ extern char &GCNRewritePartialRegUsesID;
void initializeAMDGPUWaitSGPRHazardsLegacyPass(PassRegistry &);
extern char &AMDGPUWaitSGPRHazardsLegacyID;
+class AMDGPUEliminateAGPRToVGPRCopyPass
+ : public PassInfoMixin<AMDGPUEliminateAGPRToVGPRCopyPass> {
+public:
+ AMDGPUEliminateAGPRToVGPRCopyPass() = default;
+ PreservedAnalyses run(MachineFunction &MF,
+ MachineFunctionAnalysisManager &MFAM);
+};
+
+void initializeAMDGPUEliminateAGPRToVGPRCopyLegacyPass(PassRegistry &);
+extern char &AMDGPUEliminateAGPRToVGPRCopyLegacyID;
+
class AMDGPURewriteAGPRCopyMFMAPass
: public PassInfoMixin<AMDGPURewriteAGPRCopyMFMAPass> {
public:
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUEliminateAGPRToVGPRCopy.cpp b/llvm/lib/Target/AMDGPU/AMDGPUEliminateAGPRToVGPRCopy.cpp
new file mode 100644
index 0000000000000..66e00e13149de
--- /dev/null
+++ b/llvm/lib/Target/AMDGPU/AMDGPUEliminateAGPRToVGPRCopy.cpp
@@ -0,0 +1,246 @@
+//===-- AMDGPUEliminateAGPRToVGPRCopy.cpp ---------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+/// \file \brief TODO
+///
+//===----------------------------------------------------------------------===//
+
+#include "AMDGPU.h"
+#include "GCNSubtarget.h"
+#include "SIMachineFunctionInfo.h"
+#include "SIRegisterInfo.h"
+#include "Utils/AMDGPUBaseInfo.h"
+#include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/Statistic.h"
+#include "llvm/CodeGen/LiveIntervals.h"
+#include "llvm/CodeGen/LiveRegMatrix.h"
+#include "llvm/CodeGen/MachineFunctionPass.h"
+#include "llvm/CodeGen/TargetRegisterInfo.h"
+#include "llvm/CodeGen/VirtRegMap.h"
+#include "llvm/InitializePasses.h"
+
+using namespace llvm;
+
+#define DEBUG_TYPE "amdgpu-eliminate-agpr-to-vgpr-copy"
+
+STATISTIC(NumEliminated, "Number of copies eliminated");
+
+namespace {
+
+class AMDGPUEliminateAGPRToVGPRCopyImpl {
+ const GCNSubtarget &ST;
+ const SIInstrInfo &TII;
+ const SIRegisterInfo &TRI;
+ MachineRegisterInfo &MRI;
+ VirtRegMap &VRM;
+ LiveRegMatrix &LRM;
+ LiveIntervals &LIS;
+
+public:
+ AMDGPUEliminateAGPRToVGPRCopyImpl(MachineFunction &MF, VirtRegMap &VRM,
+ LiveRegMatrix &LRM, LiveIntervals &LIS)
+ : ST(MF.getSubtarget<GCNSubtarget>()), TII(*ST.getInstrInfo()),
+ TRI(*ST.getRegisterInfo()), MRI(MF.getRegInfo()), VRM(VRM), LRM(LRM),
+ LIS(LIS) {}
+
+ bool areAllUsesCompatible(Register Reg) const;
+
+ bool run(MachineFunction &MF) const;
+};
+
+bool AMDGPUEliminateAGPRToVGPRCopyImpl::areAllUsesCompatible(
+ Register Reg) const {
+ return all_of(MRI.use_operands(Reg), [&](const MachineOperand &MO) {
+ const MachineInstr &ParentMI = *MO.getParent();
+ if (!SIInstrInfo::isMFMA(ParentMI))
+ return false;
+ return &MO == TII.getNamedOperand(ParentMI, AMDGPU::OpName::src0) ||
+ &MO == TII.getNamedOperand(ParentMI, AMDGPU::OpName::src1);
+ });
+}
+
+bool AMDGPUEliminateAGPRToVGPRCopyImpl::run(MachineFunction &MF) const {
+ // This only applies on subtargets that have a configurable AGPR vs. VGPR
+ // allocation.
+ if (!ST.hasGFX90AInsts())
+ return false;
+
+ // Early exit if no AGPRs were assigned.
+ if (!LRM.isPhysRegUsed(AMDGPU::AGPR0))
+ return false;
+
+ bool MadeChange = false;
+
+ for (MachineBasicBlock &MBB : MF) {
+ for (MachineInstr &CopyMI : make_early_inc_range(MBB)) {
+ // Find full copies...
+ if (!CopyMI.isFullCopy())
+ continue;
+
+ // ... whose destination was mapped to a VGPR or AGPR...
+ Register DstReg = CopyMI.getOperand(0).getReg();
+ if (!DstReg.isVirtual())
+ continue;
+ Register DstPhysReg = VRM.getPhys(DstReg);
+ if (!DstPhysReg)
+ continue;
+ const TargetRegisterClass *DstRC = TRI.getPhysRegBaseClass(DstPhysReg);
+ if (!TRI.hasVectorRegisters(DstRC) || TRI.hasSGPRs(DstRC))
+ continue;
+
+ // ... and whose source was mapped to an AGPR.
+ Register SrcReg = CopyMI.getOperand(1).getReg();
+ if (!SrcReg.isVirtual() || SrcReg == DstReg)
+ continue;
+ Register SrcPhysReg = VRM.getPhys(SrcReg);
+ if (!SrcPhysReg)
+ continue;
+ const TargetRegisterClass *SrcRC = TRI.getPhysRegBaseClass(SrcPhysReg);
+ if (!TRI.isAGPRClass(SrcRC))
+ continue;
+
+ bool DstIsAGPR = TRI.hasAGPRs(DstRC);
+
+ LLVM_DEBUG({
+ dbgs() << "AGPR->AVGPR copy: " << CopyMI;
+ dbgs() << " "
+ << printReg(DstReg, &TRI, CopyMI.getOperand(0).getSubReg(), &MRI)
+ << " <-> " << printReg(DstPhysReg, &TRI, 0, &MRI) << "\n";
+ dbgs() << " "
+ << printReg(SrcReg, &TRI, CopyMI.getOperand(1).getSubReg(), &MRI)
+ << " <-> " << printReg(SrcPhysReg, &TRI, 0, &MRI) << "\n";
+ });
+
+ LiveInterval &SrcLI = LIS.getInterval(SrcReg);
+ const VNInfo *SrcVNI = SrcLI.getVNInfoAt(LIS.getInstructionIndex(CopyMI));
+ assert(SrcVNI && "VNI must exist");
+
+ bool AllUsesCompatible =
+ all_of(MRI.use_operands(DstReg), [&](const MachineOperand &MO) {
+ // Destination's use must be src0/src1 operands of an MFMA or
+ // another copy.
+ const MachineInstr &UseMI = *MO.getParent();
+ if (!DstIsAGPR) {
+ if (SIInstrInfo::isMFMA(UseMI)) {
+ if (&MO != TII.getNamedOperand(UseMI, AMDGPU::OpName::src0) &&
+ &MO != TII.getNamedOperand(UseMI, AMDGPU::OpName::src1)) {
+ LLVM_DEBUG(dbgs() << " Incompatible MFMA operand: " << UseMI);
+ return false;
+ }
+ } else if (!UseMI.isFullCopy()){
+ LLVM_DEBUG(dbgs() << " Incompatible user: " << UseMI);
+ return false;
+ }
+ } else {
+ LLVM_DEBUG(dbgs() << " Skipping user check (dst is AGPR)\n");
+ }
+
+ // Source must be available at use point.
+ const VNInfo *UseVNI =
+ SrcLI.getVNInfoAt(LIS.getInstructionIndex(UseMI));
+ if (SrcVNI != UseVNI) {
+ LLVM_DEBUG(dbgs() << " AGPR no longer available at " << UseMI);
+ }
+ return true;
+ });
+ if (!AllUsesCompatible)
+ continue;
+
+ LLVM_DEBUG(dbgs() << " -> Eliminated\n");
+ ++NumEliminated;
+
+ // Remove the copy's destination register.
+ MRI.replaceRegWith(DstReg, SrcReg);
+ LRM.unassign(LIS.getInterval(DstReg));
+ LIS.removeInterval(DstReg);
+
+ // Delete the copy instruction.
+ LIS.RemoveMachineInstrFromMaps(CopyMI);
+ CopyMI.eraseFromParent();
+
+ // Recompute the source register's interval.
+ // TODO: necessary? It is already live at all uses by construction.
+ LIS.removeInterval(SrcReg);
+ LIS.createAndComputeVirtRegInterval(SrcReg);
+ MadeChange = true;
+ }
+ }
+
+ return MadeChange;
+}
+
+class AMDGPUEliminateAGPRToVGPRCopyLegacy : public MachineFunctionPass {
+public:
+ static char ID;
+
+ AMDGPUEliminateAGPRToVGPRCopyLegacy() : MachineFunctionPass(ID) {
+ initializeAMDGPUEliminateAGPRToVGPRCopyLegacyPass(
+ *PassRegistry::getPassRegistry());
+ }
+
+ bool runOnMachineFunction(MachineFunction &MF) override;
+
+ StringRef getPassName() const override {
+ return "AMDGPU Eliminate AGPR-to-VGPR Copy";
+ }
+
+ void getAnalysisUsage(AnalysisUsage &AU) const override {
+ AU.addRequired<LiveIntervalsWrapperPass>();
+ AU.addRequired<VirtRegMapWrapperLegacy>();
+ AU.addRequired<LiveRegMatrixWrapperLegacy>();
+
+ AU.addPreserved<LiveIntervalsWrapperPass>();
+ AU.addPreserved<VirtRegMapWrapperLegacy>();
+ AU.addPreserved<LiveRegMatrixWrapperLegacy>();
+ AU.setPreservesAll();
+ MachineFunctionPass::getAnalysisUsage(AU);
+ }
+};
+
+} // End anonymous namespace.
+
+INITIALIZE_PASS_BEGIN(AMDGPUEliminateAGPRToVGPRCopyLegacy, DEBUG_TYPE,
+ "AMDGPU Eliminate AGPR-to-VGPR Copy", false, false)
+INITIALIZE_PASS_DEPENDENCY(LiveIntervalsWrapperPass)
+INITIALIZE_PASS_DEPENDENCY(VirtRegMapWrapperLegacy)
+INITIALIZE_PASS_DEPENDENCY(LiveRegMatrixWrapperLegacy)
+INITIALIZE_PASS_END(AMDGPUEliminateAGPRToVGPRCopyLegacy, DEBUG_TYPE,
+ "AMDGPU Eliminate AGPR-to-VGPR Copy", false, false)
+
+char AMDGPUEliminateAGPRToVGPRCopyLegacy::ID = 0;
+
+char &llvm::AMDGPUEliminateAGPRToVGPRCopyLegacyID =
+ AMDGPUEliminateAGPRToVGPRCopyLegacy::ID;
+
+bool AMDGPUEliminateAGPRToVGPRCopyLegacy::runOnMachineFunction(
+ MachineFunction &MF) {
+ if (skipFunction(MF.getFunction()))
+ return false;
+
+ auto &VRM = getAnalysis<VirtRegMapWrapperLegacy>().getVRM();
+ auto &LRM = getAnalysis<LiveRegMatrixWrapperLegacy>().getLRM();
+ auto &LIS = getAnalysis<LiveIntervalsWrapperPass>().getLIS();
+
+ AMDGPUEliminateAGPRToVGPRCopyImpl Impl(MF, VRM, LRM, LIS);
+ return Impl.run(MF);
+}
+
+PreservedAnalyses
+AMDGPUEliminateAGPRToVGPRCopyPass::run(MachineFunction &MF,
+ MachineFunctionAnalysisManager &MFAM) {
+ VirtRegMap &VRM = MFAM.getResult<VirtRegMapAnalysis>(MF);
+ LiveRegMatrix &LRM = MFAM.getResult<LiveRegMatrixAnalysis>(MF);
+ LiveIntervals &LIS = MFAM.getResult<LiveIntervalsAnalysis>(MF);
+
+ AMDGPUEliminateAGPRToVGPRCopyImpl Impl(MF, VRM, LRM, LIS);
+ if (!Impl.run(MF))
+ return PreservedAnalyses::all();
+ auto PA = getMachineFunctionPassPreservedAnalyses();
+ PA.preserveSet<CFGAnalyses>();
+ return PA;
+}
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
index b6c6d927d0e89..34963826d16a5 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
+++ b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
@@ -106,6 +106,7 @@ MACHINE_FUNCTION_ANALYSIS("amdgpu-resource-usage", AMDGPUResourceUsageAnalysis(*
#endif
MACHINE_FUNCTION_PASS("amdgpu-insert-delay-alu", AMDGPUInsertDelayAluPass())
MACHINE_FUNCTION_PASS("amdgpu-isel", AMDGPUISelDAGToDAGPass(*this))
+MACHINE_FUNCTION_PASS("amdgpu-eliminate-agpr-to-vgpr-copy", AMDGPUEliminateAGPRToVGPRCopyPass())
MACHINE_FUNCTION_PASS("amdgpu-mark-last-scratch-load", AMDGPUMarkLastScratchLoadPass())
MACHINE_FUNCTION_PASS("amdgpu-pre-ra-long-branch-reg", GCNPreRALongBranchRegPass())
MACHINE_FUNCTION_PASS("amdgpu-reserve-wwm-regs", AMDGPUReserveWWMRegsPass())
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
index c1f17033d04a8..5512b15f4cb31 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -528,6 +528,7 @@ extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializeAMDGPUTarget() {
initializeAMDGPULowerKernelArgumentsPass(*PR);
initializeAMDGPUPromoteKernelArgumentsPass(*PR);
initializeAMDGPULowerKernelAttributesPass(*PR);
+ initializeAMDGPUEliminateAGPRToVGPRCopyLegacyPass(*PR);
initializeAMDGPUExportKernelRuntimeHandlesLegacyPass(*PR);
initializeAMDGPUPostLegalizerCombinerPass(*PR);
initializeAMDGPUPreLegalizerCombinerPass(*PR);
@@ -1594,6 +1595,7 @@ bool GCNPassConfig::addPreRewrite() {
if (EnableRegReassign)
addPass(&GCNNSAReassignID);
+ addPass(&AMDGPUEliminateAGPRToVGPRCopyLegacyID);
addPass(&AMDGPURewriteAGPRCopyMFMALegacyID);
return true;
}
diff --git a/llvm/lib/Target/AMDGPU/CMakeLists.txt b/llvm/lib/Target/AMDGPU/CMakeLists.txt
index c466f9cf0f359..abc8020920e7d 100644
--- a/llvm/lib/Target/AMDGPU/CMakeLists.txt
+++ b/llvm/lib/Target/AMDGPU/CMakeLists.txt
@@ -53,6 +53,7 @@ add_llvm_target(AMDGPUCodeGen
AMDGPUCodeGenPrepare.cpp
AMDGPUCombinerHelper.cpp
AMDGPUCtorDtorLowering.cpp
+ AMDGPUEliminateAGPRToVGPRCopy.cpp
AMDGPUExportClustering.cpp
AMDGPUExportKernelRuntimeHandles.cpp
AMDGPUFrameLowering.cpp
diff --git a/llvm/test/CodeGen/AMDGPU/llc-pipeline.ll b/llvm/test/CodeGen/AMDGPU/llc-pipeline.ll
index 2a5c65278f7dc..ed865dc62592b 100644
--- a/llvm/test/CodeGen/AMDGPU/llc-pipeline.ll
+++ b/llvm/test/CodeGen/AMDGPU/llc-pipeline.ll
@@ -377,6 +377,7 @@
; GCN-O1-NEXT: Live Register Matrix
; GCN-O1-NEXT: Greedy Register Allocator
; GCN-O1-NEXT: GCN NSA Reassign
+; GCN-O1-NEXT: AMDGPU Eliminate AGPR-to-VGPR Copy
; GCN-O1-NEXT: AMDGPU Rewrite AGPR-Copy-MFMA
; GCN-O1-NEXT: Virtual Register Rewriter
; GCN-O1-NEXT: AMDGPU Mark Last Scratch Load
@@ -689,6 +690,7 @@
; GCN-O1-OPTS-NEXT: Live Register Matrix
; GCN-O1-OPTS-NEXT: Greedy Register Allocator
; GCN-O1-OPTS-NEXT: GCN NSA Reassign
+; GCN-O1-OPTS-NEXT: AMDGPU Eliminate AGPR-to-VGPR Copy
; GCN-O1-OPTS-NEXT: AMDGPU Rewrite AGPR-Copy-MFMA
; GCN-O1-OPTS-NEXT: Virtual Register Rewriter
; GCN-O1-OPTS-NEXT: AMDGPU Mark Last Scratch Load
@@ -1007,6 +1009,7 @@
; GCN-O2-NEXT: Live Register Matrix
; GCN-O2-NEXT: Greedy Register Allocator
; GCN-O2-NEXT: GCN NSA Reassign
+; GCN-O2-NEXT: AMDGPU Eliminate AGPR-to-VGPR Copy
; GCN-O2-NEXT: AMDGPU Rewrite AGPR-Copy-MFMA
; GCN-O2-NEXT: Virtual Register Rewriter
; GCN-O2-NEXT: AMDGPU Mark Last Scratch Load
@@ -1338,6 +1341,7 @@
; GCN-O3-NEXT: Live Register Matrix
; GCN-O3-NEXT: Greedy Register Allocator
; GCN-O3-NEXT: GCN NSA Reassign
+; GCN-O3-NEXT: AMDGPU Eliminate AGPR-to-VGPR Copy
; GCN-O3-NEXT: AMDGPU Rewrite AGPR-Copy-MFMA
; GCN-O3-NEXT: Virtual Register Rewriter
; GCN-O3-NEXT: AMDGPU Mark Last Scratch Load
diff --git a/llvm/test/CodeGen/AMDGPU/mfma-loop.ll b/llvm/test/CodeGen/AMDGPU/mfma-loop.ll
index 6110b3101020a..fdf4dfd1eaefa 100644
--- a/llvm/test/CodeGen/AMDGPU/mfma-loop.ll
+++ b/llvm/test/CodeGen/AMDGPU/mfma-loop.ll
@@ -103,6 +103,9 @@ define amdgpu_kernel void @test_mfma_loop_zeroinit(ptr addrspace(1) %arg) #0 {
; GFX90A-LABEL: test_mfma_loop_zeroinit:
; GFX90A: ; %bb.0: ; %entry
; GFX90A-NEXT: v_accvgpr_write_b32 a31, 0
+; GFX90A-NEXT: s_mov_b32 s0, 16
+; GFX90A-NEXT: v_mov_b32_e32 v0, 2.0
+; GFX90A-NEXT: v_mov_b32_e32 v1, 1.0
; GFX90A-NEXT: v_accvgpr_write_b32 a30, 0
; GFX90A-NEXT: v_accvgpr_write_b32 a29, 0
; GFX90A-NEXT: v_accvgpr_write_b32 a28, 0
@@ -134,14 +137,42 @@ define amdgpu_kernel void @test_mfma_loop_zeroinit(ptr addrspace(1) %arg) #0 {
; GFX90A-NEXT: v_accvgpr_write_b32 a2, 0
; GFX90A-NEXT: v_accvgpr_write_b32 a1, 0
; GFX90A-NEXT: v_accvgpr_write_b32 a0, 0
-; GFX90A-NEXT: s_mov_b32 s0, 16
-; GFX90A-NEXT: v_mov_b32_e32 v0, 2.0
-; GFX90A-NEXT: v_mov_b32_e32 v1, 1.0
; GFX90A-NEXT: .LBB0_1: ; %for.cond.preheader
; GFX90A-NEXT: ; =>This Inner Loop Header: Depth=1
-; GFX90A-NEXT: s_nop 1
-; GFX90A-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v0, a[0:31]
+; GFX90A-NEXT: ; kill: def $agpr30 killed $agpr30
+; GFX90A-NEXT: ; kill: def $agpr29 killed $agpr29
+; GFX90A-NEXT: ; kill: def $agpr28 killed $agpr28
+; GFX90A-NEXT: ; kill: def $agpr27 killed $agpr27
+; GFX90A-NEXT: ; kill: def $agpr26 killed $agpr26
+; GFX90A-NEXT: ; kill: def $agpr25 killed $agpr25
+; GFX90A-NEXT: ; kill: def $agpr24 killed $agpr24
+; GFX90A-NEXT: ; kill: def $agpr23 killed $agpr23
+; GFX90A-NEXT: ; kill: def $agpr22 killed $agpr22
+; GFX90A-NEXT: ; kill: def $agpr21 killed $agpr21
+; GFX90A-NEXT: ; kill: def $agpr20 killed $agpr20
+; GFX90A-NEXT: ; kill: def $agpr19 killed $agpr19
+; GFX90A-NEXT: ; kill: def $agpr18 killed $agpr18
+; GFX90A-NEXT: ; kill: def $agpr17 killed $agpr17
+; GFX90A-NEXT: ; kill: def $agpr16 killed $agpr16
+; GFX90A-NEXT: ; kill: def $agpr15 killed $agpr15
+; GFX90A-NEXT: ; kill: def $agpr14 killed $agpr14
+; GFX90A-NEXT: ; kill: def $agpr13 killed $agpr13
+; GFX90A-NEXT: ; kill: def $agpr12 killed $agpr12
+; GFX90A-NEXT: ; kill: def $agpr11 killed $agpr11
+; GFX90A-NEXT: ; kill: def $agpr10 killed $agpr10
+; GFX90A-NEXT: ; kill: def $agpr9 killed $agpr9
+; GFX90A-NEXT: ; kill: def $agpr8 killed $agpr8
+; GFX90A-NEXT: ; kill: def $agpr7 killed $agpr7
+; GFX90A-NEXT: ; kill: def $agpr6 killed $agpr6
+; GFX90A-NEXT: ; kill: def $agpr5 killed $agpr5
+; GFX90A-NEXT: ; kill: def $agpr4 killed $agpr4
+; GFX90A-NEXT: ; kill: def $agpr3 killed $agpr3
+; GFX90A-NEXT: ; kill: def $agpr2 killed $agpr2
+; GFX90A-NEXT: ; kill: def $agpr1 killed $agpr1
+; GFX90A-NEXT: ; kill: def $agpr0 killed $agpr0
; GFX90A-NEXT: s_add_i32 s0, s0, -1
+; GFX90A-NEXT: s_nop 0
+; GFX90A-NEXT: v_mfma_f32_32x32x1f32 a[0:31], v1, v0, a[0:31]
; GFX90A-NEXT: s_cmp_lg_u32 s0, 0
; GFX90A-NEXT: s_cbranch_scc1 .LBB0_1
; GFX90A-NEXT: ; %bb.2: ; %exit
@@ -149,7 +180,7 @@ define amdgpu_kernel void @test_mfma_loop_zeroinit(ptr addrspace(1) %arg) #0 {
; GFX90A-NEXT: v_mov_b32_e32 v0, 0
; GFX90A-NEXT: s_waitcnt lgkmcnt(0)
; GFX90A-NEXT: s_nop 7
-; GFX90A-NEXT: s_nop 4
+; GFX90A-NEXT: s_nop 5
; GFX90A-NEXT: global_store_dwordx4 v0, a[28:31], s[0:1] offset:112
; GFX90A-NEXT: global_store_dwordx4 v0, a[24:27], s[0:1] offset:96
; GFX90A-NEXT: global_store_dwordx4 v0, a[20:23], s[0:1] offset:80
@@ -163,6 +194,9 @@ define amdgpu_kernel void @test_mfma_loop_zeroinit(ptr addrspace(1) %arg) #0 {
; GFX942-LABEL: test_mfma_loop_zeroinit:
; GFX942: ; %bb.0: ; %entry
; GFX942-NEXT: v_accvgpr_write_b32 a31, 0
+; GFX942-NEXT: s_mov_b32 s0, 16
+; GFX942-NEXT: v_mov_b32_e32 v0, 2.0
+; GFX942-NEXT: v_mov_b32_e32 v1, 1.0
; GFX942-NEXT: v_accvgpr_write_b32 a30, 0
; GFX942-NEXT: v_accvgpr_write_b32 a29, 0
; GFX942-NEXT: v_accvgpr_write_b32 a28, 0
@@ -194,14 +228,42 @@ define amdgpu_kernel void @test_mfma_loop_zeroinit(ptr addrspace(1) %arg) #0 {
; GFX942-NEXT: v_accvgpr_write_b32 a2, 0
; GFX942-NEXT: v_accvgpr_write_b32 a1, 0
; GFX942-NEXT: v_accvgpr_write_b32 a0, 0
-; GFX942-NEXT: s_mov_b32 s0, 16
-; GFX942-NEXT: v_mov_b32_e32 v0, 2.0
-; GFX942-NEXT: v_mov_b32_e32 v1, 1.0
; GFX942-NEXT: .LBB0_1: ; %for.cond.preheader
; GFX942-NEXT: ; =>This Inner Loop Header: Depth=1
-; GFX942-NEXT: s_nop 1
-; GFX942-NEXT: v_mfma_f32_32x32x1_2b_f32 a[0:31], v1, v0, a[0:31]
+; GFX942-NEXT: ; kill: def $agpr30 killed $agpr30
+; GFX942-NEXT: ; kill: def $agpr29 killed $agpr29
+; GFX942-NEXT: ; kill: def $agpr28 killed $agpr28
+; GFX942-NEXT: ; kill: def $agpr27 killed $agpr27
+; GFX942-NEXT: ; kill: def $agpr26 killed $agpr26
+; GFX942-NEXT: ; kill: def $agpr25 killed $agpr25
+; GFX942-NEXT: ; kill: def $agpr24 killed $agpr24
+; GFX942-NEXT: ; kill: def $agpr23 killed $agpr23
+; GFX942-NEXT: ; kill: def $agpr22 killed $agpr22
+; GFX942-NEXT: ; kill: def $agpr21 killed $agpr21
+; GFX942-NEXT: ; kill: def $agpr20 killed $agpr20
+; GFX942-NEXT: ; kill: def $agpr19 killed $agpr19
+; GFX942-NEXT: ; kill: def $agpr18 killed $agpr18
+; GFX942-NEXT: ; kill: def $agpr17 killed $agpr17
+; GFX942-NEXT: ; kill: def $agpr16 killed $agpr16
+; GFX942-NEXT: ; kill: def $agpr15 killed $agpr15
+; GFX942-NEXT: ; kill: def $agpr14 killed $agpr14
+; GFX942-NEXT: ; kill: def $agpr13 killed $agpr13
+; GFX942-NEXT: ; kill: def $agpr12 killed $agpr12
+; GFX942-NEXT: ; kill: def $agpr11 killed $agpr11
+; GFX942-NEXT: ; kill: def $agpr10 killed $agpr10
+; GFX942-NEXT: ; kill: def $agpr9 killed $agpr9
+; GFX942-NEXT: ; kill: def $agpr8 killed $agpr8
+; GFX942-NEXT: ; kill: def $agpr7 killed $agpr7
+; GFX942-NEXT: ; kill: def $agpr6 killed $agpr6
+; GFX942-NEXT: ; kill: def $agpr5 killed $agpr5
+; GFX942-NEXT: ; kill: def $agpr4 killed $agpr4
+; GFX942-NEXT: ; kill: def $agpr3 killed $agpr3
+; GFX942-NEXT: ; kill: def $agpr2 killed $agpr2
+; GFX942-NEXT: ; kill: def $agpr1 killed $agpr1
+; GFX942-NEXT: ; kill: def $agpr0 killed $agpr0
; GFX942-NEXT: s_add_i32 s0, s0, -1
+; GFX942-NEXT: s_nop 0
+; GFX942-NEXT: v_mfma_f32_32x32x1_2b_f32 a[0:31], v1...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/153901
More information about the llvm-commits
mailing list