[llvm] [AMDGPU] RewriteMFMAFormStage: single exit copy for same-block uses (PR #209179)
Romanov Vlad via llvm-commits
llvm-commits at lists.llvm.org
Tue Jul 21 03:21:37 PDT 2026
https://github.com/romanovvlad updated https://github.com/llvm/llvm-project/pull/209179
>From cb0ced7e21910664eb2544b9ab8c25ead1f97d58 Mon Sep 17 00:00:00 2001
From: Vlad <Vladislav.Romanov at amd.com>
Date: Mon, 13 Jul 2026 04:29:03 -0500
Subject: [PATCH 1/4] [AMDGPU] RewriteMFMAFormStage: share single exit copy for
same-block uses
When multiple VGPR-requiring uses of an MFMA dst are in the same block
as the MFMA, share a single AGPR->VGPR copy placed before the earliest
use, instead of creating one per use.
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 39 ++++++---
.../AMDGPU/mfma-no-multiple-copies.mir | 84 +++++++++++++++++++
2 files changed, 111 insertions(+), 12 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index a4f854beaeebe..d6189f8875dab 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -2810,6 +2810,9 @@ bool RewriteMFMAFormStage::rewrite(
}
DenseSet<MachineOperand *> &DstRegSet = ReplaceMap[DstReg];
+ // One AGPR→VGPR copy per dst register, shared by all same-block uses.
+ Register SameBlockCopyReg;
+ MachineInstr *EarliestSameBlockUse = nullptr;
for (MachineOperand *RU : DstReachingUseCopies) {
MachineBasicBlock *RUBlock = RU->getParent()->getParent();
// Just keep track of the reaching use of this register by block. After we
@@ -2819,22 +2822,34 @@ bool RewriteMFMAFormStage::rewrite(
continue;
}
- // Special case, the use is in the same block as the MFMA. Insert the copy
- // just before the use.
- const TargetRegisterClass *DstRC = DAG.MRI.getRegClass(DstReg);
- const TargetRegisterClass *VGPRRC = SRI->getEquivalentVGPRClass(DstRC);
- Register NewUseReg = DAG.MRI.createVirtualRegister(VGPRRC);
+ // Lazily create the copy register on first same-block use.
+ if (!SameBlockCopyReg.isValid()) {
+ const TargetRegisterClass *DstRC = DAG.MRI.getRegClass(DstReg);
+ const TargetRegisterClass *VGPRRC = SRI->getEquivalentVGPRClass(DstRC);
+ SameBlockCopyReg = DAG.MRI.createVirtualRegister(VGPRRC);
+ }
+
+ // Track the earliest use for copy insertion point.
MachineInstr *UseInst = RU->getParent();
+ if (!EarliestSameBlockUse ||
+ SlotIndex::isEarlierInstr(
+ DAG.LIS->getInstructionIndex(*UseInst),
+ DAG.LIS->getInstructionIndex(*EarliestSameBlockUse)))
+ EarliestSameBlockUse = UseInst;
+
+ RU->setReg(SameBlockCopyReg);
+ }
+
+ // Insert the copy before the earliest same-block use.
+ if (SameBlockCopyReg.isValid()) {
MachineInstrBuilder VGPRCopy =
- BuildMI(*UseInst->getParent(), UseInst->getIterator(),
- UseInst->getDebugLoc(), TII->get(TargetOpcode::COPY))
- .addDef(NewUseReg, {}, 0)
+ BuildMI(*EarliestSameBlockUse->getParent(),
+ EarliestSameBlockUse->getIterator(),
+ EarliestSameBlockUse->getDebugLoc(),
+ TII->get(TargetOpcode::COPY))
+ .addDef(SameBlockCopyReg, {}, 0)
.addUse(DstReg, {}, 0);
DAG.LIS->InsertMachineInstrInMaps(*VGPRCopy);
- // Since we know this use has only one reaching def, we can replace the
- // use reg.
- RU->setReg(NewUseReg);
- // Track the copy source operand for r eplacement.
DstRegSet.insert(&VGPRCopy->getOperand(1));
}
diff --git a/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir b/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir
new file mode 100644
index 0000000000000..98ad1633cbfc5
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir
@@ -0,0 +1,84 @@
+# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 6
+# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -run-pass=machine-scheduler -amdgpu-disable-rewrite-mfma-form-sched-stage=false %s -o - | FileCheck %s
+
+# Test that only one copy generated for all users of mfma in the same block.
+
+--- |
+ define void @mfma_no_multiple_copies() #0 { ret void }
+ attributes #0 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" }
+...
+
+---
+name: mfma_no_multiple_copies
+tracksRegLiveness: true
+machineFunctionInfo:
+ isEntryFunction: true
+body: |
+ ; CHECK-LABEL: name: mfma_no_multiple_copies
+ ; CHECK: bb.0:
+ ; CHECK-NEXT: successors: %bb.1(0x80000000)
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: %p0:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %m3:vreg_128_align2 = IMPLICIT_DEF
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = COPY %m3
+ ; CHECK-NEXT: %p1:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p2:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p3:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p4:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p5:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p6:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: %p7:vreg_1024 = IMPLICIT_DEF
+ ; CHECK-NEXT: S_BRANCH %bb.1
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: bb.1:
+ ; CHECK-NEXT: successors: %bb.1(0x7c000000), %bb.2(0x04000000)
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: %m0:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, [[COPY]], 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %m1:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m0, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %m2:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m1, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: S_CBRANCH_SCC1 %bb.2, implicit undef $scc
+ ; CHECK-NEXT: S_BRANCH %bb.1
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: bb.2:
+ ; CHECK-NEXT: KILL %p0, %p1, %p2, %p3, %p4, %p5, %p6, %p7
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m2, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY [[COPY]]
+ ; CHECK-NEXT: %out1:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, [[COPY1]], 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %out2:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s1:vreg_128_align2, undef %s0:vreg_128_align2, [[COPY1]], 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: S_ENDPGM 0, implicit %out1, implicit %out2
+ bb.0:
+ successors: %bb.1(0x80000000)
+ %p0:vreg_1024 = IMPLICIT_DEF
+ %p1:vreg_1024 = IMPLICIT_DEF
+ %p2:vreg_1024 = IMPLICIT_DEF
+ %p3:vreg_1024 = IMPLICIT_DEF
+ %p4:vreg_1024 = IMPLICIT_DEF
+ %p5:vreg_1024 = IMPLICIT_DEF
+ %p6:vreg_1024 = IMPLICIT_DEF
+ %p7:vreg_1024 = IMPLICIT_DEF
+
+ %m3:vreg_128_align2 = IMPLICIT_DEF
+
+ S_BRANCH %bb.1
+
+ bb.1:
+ successors: %bb.1(0x7c000000), %bb.2(0x04000000)
+
+ %m0:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m3, 0, 0, 0, implicit $mode, implicit $exec
+ %m1:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m0, 0, 0, 0, implicit $mode, implicit $exec
+ %m2:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m1, 0, 0, 0, implicit $mode, implicit $exec
+
+ S_CBRANCH_SCC1 %bb.2, implicit undef $scc
+ S_BRANCH %bb.1
+
+ bb.2:
+
+ %m3:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m2, 0, 0, 0, implicit $mode, implicit $exec
+
+ %out1:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m3, 0, 0, 0, implicit $mode, implicit $exec
+ %out2:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s1:vreg_128_align2, undef %s0:vreg_128_align2, %m3, 0, 0, 0, implicit $mode, implicit $exec
+
+ KILL %p0, %p1, %p2, %p3, %p4, %p5, %p6, %p7
+
+ S_ENDPGM 0, implicit %out1, implicit %out2
+...
>From b9663e51adae0ae512adfc1689dbf5d66aa99e9b Mon Sep 17 00:00:00 2001
From: Vlad <Vladislav.Romanov at amd.com>
Date: Fri, 17 Jul 2026 05:19:39 -0500
Subject: [PATCH 2/4] improve
---
llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp | 3 +--
1 file changed, 1 insertion(+), 2 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
index d6189f8875dab..b3c6b0aa2ebac 100644
--- a/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNSchedStrategy.cpp
@@ -2846,8 +2846,7 @@ bool RewriteMFMAFormStage::rewrite(
BuildMI(*EarliestSameBlockUse->getParent(),
EarliestSameBlockUse->getIterator(),
EarliestSameBlockUse->getDebugLoc(),
- TII->get(TargetOpcode::COPY))
- .addDef(SameBlockCopyReg, {}, 0)
+ TII->get(TargetOpcode::COPY), SameBlockCopyReg)
.addUse(DstReg, {}, 0);
DAG.LIS->InsertMachineInstrInMaps(*VGPRCopy);
DstRegSet.insert(&VGPRCopy->getOperand(1));
>From c41302e588122742c4d235f4f8f50bbc5339be08 Mon Sep 17 00:00:00 2001
From: Vlad <Vladislav.Romanov at amd.com>
Date: Fri, 17 Jul 2026 05:25:45 -0500
Subject: [PATCH 3/4] update test
---
llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir | 12 ++++++------
1 file changed, 6 insertions(+), 6 deletions(-)
diff --git a/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir b/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir
index 98ad1633cbfc5..8a2e3e3ece974 100644
--- a/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir
+++ b/llvm/test/CodeGen/AMDGPU/mfma-no-multiple-copies.mir
@@ -4,7 +4,7 @@
# Test that only one copy generated for all users of mfma in the same block.
--- |
- define void @mfma_no_multiple_copies() #0 { ret void }
+ define amdgpu_kernel void @mfma_no_multiple_copies() #0 { ret void }
attributes #0 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" }
...
@@ -41,8 +41,8 @@ body: |
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: KILL %p0, %p1, %p2, %p3, %p4, %p5, %p6, %p7
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m2, 0, 0, 0, implicit $mode, implicit $exec
- ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY [[COPY]]
+ ; CHECK-NEXT: %m4:areg_128_align2 = V_MFMA_F32_16X16X32_F16_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m2, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:vreg_128_align2 = COPY %m4
; CHECK-NEXT: %out1:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, [[COPY1]], 0, 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: %out2:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s1:vreg_128_align2, undef %s0:vreg_128_align2, [[COPY1]], 0, 0, 0, implicit $mode, implicit $exec
; CHECK-NEXT: S_ENDPGM 0, implicit %out1, implicit %out2
@@ -73,10 +73,10 @@ body: |
bb.2:
- %m3:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m2, 0, 0, 0, implicit $mode, implicit $exec
+ %m4:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m2, 0, 0, 0, implicit $mode, implicit $exec
- %out1:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m3, 0, 0, 0, implicit $mode, implicit $exec
- %out2:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s1:vreg_128_align2, undef %s0:vreg_128_align2, %m3, 0, 0, 0, implicit $mode, implicit $exec
+ %out1:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s0:vreg_128_align2, undef %s1:vreg_128_align2, %m4, 0, 0, 0, implicit $mode, implicit $exec
+ %out2:vreg_128_align2 = V_MFMA_F32_16X16X32_F16_vgprcd_e64 undef %s1:vreg_128_align2, undef %s0:vreg_128_align2, %m4, 0, 0, 0, implicit $mode, implicit $exec
KILL %p0, %p1, %p2, %p3, %p4, %p5, %p6, %p7
>From a5cfc7b021c5b78127b1e691b5db2fb8a90bd5d0 Mon Sep 17 00:00:00 2001
From: Vlad <Vladislav.Romanov at amd.com>
Date: Tue, 21 Jul 2026 05:15:59 -0500
Subject: [PATCH 4/4] add llvm ir test
---
.../AMDGPU/rewrite-mfma-single-exit-copy.ll | 603 ++++++++++++++++++
1 file changed, 603 insertions(+)
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-mfma-single-exit-copy.ll
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-single-exit-copy.ll b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-single-exit-copy.ll
new file mode 100644
index 0000000000000..0e2b327e4f26e
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-single-exit-copy.ll
@@ -0,0 +1,603 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -mtriple=amdgpu9.50-amd-amdhsa -amdgpu-disable-rewrite-mfma-form-sched-stage=false < %s | FileCheck %s
+
+; Test that the MFMA rewrite stage generates only one AGPR->VGPR copy
+; when multiple VGPR-requiring instructions in the exit block use the
+; same MFMA result.
+;
+; 8 x <32 x float> loop-carried carriers create 256 VGPRs of pressure.
+; 3 chained MFMAs in the loop, then in the exit block 1 MFMA (m3)
+; followed by 2 MFMAs (out1, out2) that both use m3 as src2, plus
+; fadd uses of out1/out2 that require VGPRs.
+; The two exit MFMAs should share one v_accvgpr_read copy of m3.
+
+declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half>, <8 x half>, <4 x float>, i32 immarg, i32 immarg, i32 immarg)
+
+define amdgpu_kernel void @single_exit_copy(
+; CHECK-LABEL: single_exit_copy:
+; CHECK: ; %bb.0: ; %entry
+; CHECK-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x0
+; CHECK-NEXT: s_load_dword s0, s[4:5], 0x20
+; CHECK-NEXT: s_mov_b32 s1, 0
+; CHECK-NEXT: v_mov_b32_e32 v31, 0
+; CHECK-NEXT: v_mov_b32_e32 v30, 0
+; CHECK-NEXT: v_mov_b32_e32 v29, 0
+; CHECK-NEXT: v_mov_b32_e32 v28, 0
+; CHECK-NEXT: v_mov_b32_e32 v27, 0
+; CHECK-NEXT: v_mov_b32_e32 v26, 0
+; CHECK-NEXT: v_mov_b32_e32 v25, 0
+; CHECK-NEXT: v_mov_b32_e32 v24, 0
+; CHECK-NEXT: v_mov_b32_e32 v23, 0
+; CHECK-NEXT: v_mov_b32_e32 v22, 0
+; CHECK-NEXT: v_mov_b32_e32 v21, 0
+; CHECK-NEXT: v_mov_b32_e32 v20, 0
+; CHECK-NEXT: v_mov_b32_e32 v19, 0
+; CHECK-NEXT: v_mov_b32_e32 v18, 0
+; CHECK-NEXT: v_mov_b32_e32 v17, 0
+; CHECK-NEXT: v_mov_b32_e32 v16, 0
+; CHECK-NEXT: v_mov_b32_e32 v15, 0
+; CHECK-NEXT: v_mov_b32_e32 v14, 0
+; CHECK-NEXT: v_mov_b32_e32 v13, 0
+; CHECK-NEXT: v_mov_b32_e32 v12, 0
+; CHECK-NEXT: v_mov_b32_e32 v11, 0
+; CHECK-NEXT: v_mov_b32_e32 v10, 0
+; CHECK-NEXT: v_mov_b32_e32 v9, 0
+; CHECK-NEXT: v_mov_b32_e32 v8, 0
+; CHECK-NEXT: v_mov_b32_e32 v7, 0
+; CHECK-NEXT: v_mov_b32_e32 v6, 0
+; CHECK-NEXT: v_mov_b32_e32 v5, 0
+; CHECK-NEXT: v_mov_b32_e32 v4, 0
+; CHECK-NEXT: v_mov_b32_e32 v3, 0
+; CHECK-NEXT: v_mov_b32_e32 v2, 0
+; CHECK-NEXT: v_mov_b32_e32 v1, 0
+; CHECK-NEXT: v_mov_b32_e32 v0, 0
+; CHECK-NEXT: v_mov_b32_e32 v63, 0
+; CHECK-NEXT: v_mov_b32_e32 v62, 0
+; CHECK-NEXT: v_mov_b32_e32 v61, 0
+; CHECK-NEXT: v_mov_b32_e32 v60, 0
+; CHECK-NEXT: v_mov_b32_e32 v59, 0
+; CHECK-NEXT: v_mov_b32_e32 v58, 0
+; CHECK-NEXT: v_mov_b32_e32 v57, 0
+; CHECK-NEXT: v_mov_b32_e32 v56, 0
+; CHECK-NEXT: v_mov_b32_e32 v55, 0
+; CHECK-NEXT: v_mov_b32_e32 v54, 0
+; CHECK-NEXT: v_mov_b32_e32 v53, 0
+; CHECK-NEXT: v_mov_b32_e32 v52, 0
+; CHECK-NEXT: v_mov_b32_e32 v51, 0
+; CHECK-NEXT: v_mov_b32_e32 v50, 0
+; CHECK-NEXT: v_mov_b32_e32 v49, 0
+; CHECK-NEXT: v_mov_b32_e32 v48, 0
+; CHECK-NEXT: v_mov_b32_e32 v47, 0
+; CHECK-NEXT: v_mov_b32_e32 v46, 0
+; CHECK-NEXT: v_mov_b32_e32 v45, 0
+; CHECK-NEXT: v_mov_b32_e32 v44, 0
+; CHECK-NEXT: v_mov_b32_e32 v43, 0
+; CHECK-NEXT: v_mov_b32_e32 v42, 0
+; CHECK-NEXT: v_mov_b32_e32 v41, 0
+; CHECK-NEXT: v_mov_b32_e32 v40, 0
+; CHECK-NEXT: v_mov_b32_e32 v39, 0
+; CHECK-NEXT: v_mov_b32_e32 v38, 0
+; CHECK-NEXT: v_mov_b32_e32 v37, 0
+; CHECK-NEXT: v_mov_b32_e32 v36, 0
+; CHECK-NEXT: v_mov_b32_e32 v35, 0
+; CHECK-NEXT: v_mov_b32_e32 v34, 0
+; CHECK-NEXT: v_mov_b32_e32 v33, 0
+; CHECK-NEXT: v_mov_b32_e32 v32, 0
+; CHECK-NEXT: v_mov_b32_e32 v95, 0
+; CHECK-NEXT: v_mov_b32_e32 v94, 0
+; CHECK-NEXT: v_mov_b32_e32 v93, 0
+; CHECK-NEXT: v_mov_b32_e32 v92, 0
+; CHECK-NEXT: v_mov_b32_e32 v91, 0
+; CHECK-NEXT: v_mov_b32_e32 v90, 0
+; CHECK-NEXT: v_mov_b32_e32 v89, 0
+; CHECK-NEXT: v_mov_b32_e32 v88, 0
+; CHECK-NEXT: v_mov_b32_e32 v87, 0
+; CHECK-NEXT: v_mov_b32_e32 v86, 0
+; CHECK-NEXT: v_mov_b32_e32 v85, 0
+; CHECK-NEXT: v_mov_b32_e32 v84, 0
+; CHECK-NEXT: v_mov_b32_e32 v83, 0
+; CHECK-NEXT: v_mov_b32_e32 v82, 0
+; CHECK-NEXT: v_mov_b32_e32 v81, 0
+; CHECK-NEXT: v_mov_b32_e32 v80, 0
+; CHECK-NEXT: v_mov_b32_e32 v79, 0
+; CHECK-NEXT: v_mov_b32_e32 v78, 0
+; CHECK-NEXT: v_mov_b32_e32 v77, 0
+; CHECK-NEXT: v_mov_b32_e32 v76, 0
+; CHECK-NEXT: v_mov_b32_e32 v75, 0
+; CHECK-NEXT: v_mov_b32_e32 v74, 0
+; CHECK-NEXT: v_mov_b32_e32 v73, 0
+; CHECK-NEXT: v_mov_b32_e32 v72, 0
+; CHECK-NEXT: v_mov_b32_e32 v71, 0
+; CHECK-NEXT: v_mov_b32_e32 v70, 0
+; CHECK-NEXT: v_mov_b32_e32 v69, 0
+; CHECK-NEXT: v_mov_b32_e32 v68, 0
+; CHECK-NEXT: v_mov_b32_e32 v67, 0
+; CHECK-NEXT: v_mov_b32_e32 v66, 0
+; CHECK-NEXT: v_mov_b32_e32 v65, 0
+; CHECK-NEXT: v_mov_b32_e32 v64, 0
+; CHECK-NEXT: v_mov_b32_e32 v127, 0
+; CHECK-NEXT: v_mov_b32_e32 v126, 0
+; CHECK-NEXT: v_mov_b32_e32 v125, 0
+; CHECK-NEXT: v_mov_b32_e32 v124, 0
+; CHECK-NEXT: v_mov_b32_e32 v123, 0
+; CHECK-NEXT: v_mov_b32_e32 v122, 0
+; CHECK-NEXT: v_mov_b32_e32 v121, 0
+; CHECK-NEXT: v_mov_b32_e32 v120, 0
+; CHECK-NEXT: v_mov_b32_e32 v119, 0
+; CHECK-NEXT: v_mov_b32_e32 v118, 0
+; CHECK-NEXT: v_mov_b32_e32 v117, 0
+; CHECK-NEXT: v_mov_b32_e32 v116, 0
+; CHECK-NEXT: v_mov_b32_e32 v115, 0
+; CHECK-NEXT: v_mov_b32_e32 v114, 0
+; CHECK-NEXT: v_mov_b32_e32 v113, 0
+; CHECK-NEXT: v_mov_b32_e32 v112, 0
+; CHECK-NEXT: v_mov_b32_e32 v111, 0
+; CHECK-NEXT: v_mov_b32_e32 v110, 0
+; CHECK-NEXT: v_mov_b32_e32 v109, 0
+; CHECK-NEXT: v_mov_b32_e32 v108, 0
+; CHECK-NEXT: v_mov_b32_e32 v107, 0
+; CHECK-NEXT: v_mov_b32_e32 v106, 0
+; CHECK-NEXT: v_mov_b32_e32 v105, 0
+; CHECK-NEXT: v_mov_b32_e32 v104, 0
+; CHECK-NEXT: v_mov_b32_e32 v103, 0
+; CHECK-NEXT: v_mov_b32_e32 v102, 0
+; CHECK-NEXT: v_mov_b32_e32 v101, 0
+; CHECK-NEXT: v_mov_b32_e32 v100, 0
+; CHECK-NEXT: v_mov_b32_e32 v99, 0
+; CHECK-NEXT: v_mov_b32_e32 v98, 0
+; CHECK-NEXT: v_mov_b32_e32 v97, 0
+; CHECK-NEXT: v_mov_b32_e32 v96, 0
+; CHECK-NEXT: v_mov_b32_e32 v159, 0
+; CHECK-NEXT: v_mov_b32_e32 v158, 0
+; CHECK-NEXT: v_mov_b32_e32 v157, 0
+; CHECK-NEXT: v_mov_b32_e32 v156, 0
+; CHECK-NEXT: v_mov_b32_e32 v155, 0
+; CHECK-NEXT: v_mov_b32_e32 v154, 0
+; CHECK-NEXT: v_mov_b32_e32 v153, 0
+; CHECK-NEXT: v_mov_b32_e32 v152, 0
+; CHECK-NEXT: v_mov_b32_e32 v151, 0
+; CHECK-NEXT: v_mov_b32_e32 v150, 0
+; CHECK-NEXT: v_mov_b32_e32 v149, 0
+; CHECK-NEXT: v_mov_b32_e32 v148, 0
+; CHECK-NEXT: v_mov_b32_e32 v147, 0
+; CHECK-NEXT: v_mov_b32_e32 v146, 0
+; CHECK-NEXT: v_mov_b32_e32 v145, 0
+; CHECK-NEXT: v_mov_b32_e32 v144, 0
+; CHECK-NEXT: v_mov_b32_e32 v143, 0
+; CHECK-NEXT: v_mov_b32_e32 v142, 0
+; CHECK-NEXT: v_mov_b32_e32 v141, 0
+; CHECK-NEXT: v_mov_b32_e32 v140, 0
+; CHECK-NEXT: v_mov_b32_e32 v139, 0
+; CHECK-NEXT: v_mov_b32_e32 v138, 0
+; CHECK-NEXT: v_mov_b32_e32 v137, 0
+; CHECK-NEXT: v_mov_b32_e32 v136, 0
+; CHECK-NEXT: v_mov_b32_e32 v135, 0
+; CHECK-NEXT: v_mov_b32_e32 v134, 0
+; CHECK-NEXT: v_mov_b32_e32 v133, 0
+; CHECK-NEXT: v_mov_b32_e32 v132, 0
+; CHECK-NEXT: v_mov_b32_e32 v131, 0
+; CHECK-NEXT: v_mov_b32_e32 v130, 0
+; CHECK-NEXT: v_mov_b32_e32 v129, 0
+; CHECK-NEXT: v_mov_b32_e32 v128, 0
+; CHECK-NEXT: v_mov_b32_e32 v191, 0
+; CHECK-NEXT: v_mov_b32_e32 v190, 0
+; CHECK-NEXT: v_mov_b32_e32 v189, 0
+; CHECK-NEXT: v_mov_b32_e32 v188, 0
+; CHECK-NEXT: v_mov_b32_e32 v187, 0
+; CHECK-NEXT: v_mov_b32_e32 v186, 0
+; CHECK-NEXT: v_mov_b32_e32 v185, 0
+; CHECK-NEXT: v_mov_b32_e32 v184, 0
+; CHECK-NEXT: v_mov_b32_e32 v183, 0
+; CHECK-NEXT: v_mov_b32_e32 v182, 0
+; CHECK-NEXT: v_mov_b32_e32 v181, 0
+; CHECK-NEXT: v_mov_b32_e32 v180, 0
+; CHECK-NEXT: v_mov_b32_e32 v179, 0
+; CHECK-NEXT: v_mov_b32_e32 v178, 0
+; CHECK-NEXT: v_mov_b32_e32 v177, 0
+; CHECK-NEXT: v_mov_b32_e32 v176, 0
+; CHECK-NEXT: v_mov_b32_e32 v175, 0
+; CHECK-NEXT: v_mov_b32_e32 v174, 0
+; CHECK-NEXT: v_mov_b32_e32 v173, 0
+; CHECK-NEXT: v_mov_b32_e32 v172, 0
+; CHECK-NEXT: v_mov_b32_e32 v171, 0
+; CHECK-NEXT: v_mov_b32_e32 v170, 0
+; CHECK-NEXT: v_mov_b32_e32 v169, 0
+; CHECK-NEXT: v_mov_b32_e32 v168, 0
+; CHECK-NEXT: v_mov_b32_e32 v167, 0
+; CHECK-NEXT: v_mov_b32_e32 v166, 0
+; CHECK-NEXT: v_mov_b32_e32 v165, 0
+; CHECK-NEXT: v_mov_b32_e32 v164, 0
+; CHECK-NEXT: v_mov_b32_e32 v163, 0
+; CHECK-NEXT: v_mov_b32_e32 v162, 0
+; CHECK-NEXT: v_mov_b32_e32 v161, 0
+; CHECK-NEXT: v_mov_b32_e32 v160, 0
+; CHECK-NEXT: v_mov_b32_e32 v223, 0
+; CHECK-NEXT: v_mov_b32_e32 v222, 0
+; CHECK-NEXT: v_mov_b32_e32 v221, 0
+; CHECK-NEXT: v_mov_b32_e32 v220, 0
+; CHECK-NEXT: v_mov_b32_e32 v219, 0
+; CHECK-NEXT: v_mov_b32_e32 v218, 0
+; CHECK-NEXT: v_mov_b32_e32 v217, 0
+; CHECK-NEXT: v_mov_b32_e32 v216, 0
+; CHECK-NEXT: v_mov_b32_e32 v215, 0
+; CHECK-NEXT: v_mov_b32_e32 v214, 0
+; CHECK-NEXT: v_mov_b32_e32 v213, 0
+; CHECK-NEXT: v_mov_b32_e32 v212, 0
+; CHECK-NEXT: v_mov_b32_e32 v211, 0
+; CHECK-NEXT: v_mov_b32_e32 v210, 0
+; CHECK-NEXT: v_mov_b32_e32 v209, 0
+; CHECK-NEXT: v_mov_b32_e32 v208, 0
+; CHECK-NEXT: v_mov_b32_e32 v207, 0
+; CHECK-NEXT: v_mov_b32_e32 v206, 0
+; CHECK-NEXT: v_mov_b32_e32 v205, 0
+; CHECK-NEXT: v_mov_b32_e32 v204, 0
+; CHECK-NEXT: v_mov_b32_e32 v203, 0
+; CHECK-NEXT: v_mov_b32_e32 v202, 0
+; CHECK-NEXT: v_mov_b32_e32 v201, 0
+; CHECK-NEXT: v_mov_b32_e32 v200, 0
+; CHECK-NEXT: v_mov_b32_e32 v199, 0
+; CHECK-NEXT: v_mov_b32_e32 v198, 0
+; CHECK-NEXT: v_mov_b32_e32 v197, 0
+; CHECK-NEXT: v_mov_b32_e32 v196, 0
+; CHECK-NEXT: v_mov_b32_e32 v195, 0
+; CHECK-NEXT: v_mov_b32_e32 v194, 0
+; CHECK-NEXT: v_mov_b32_e32 v193, 0
+; CHECK-NEXT: v_mov_b32_e32 v192, 0
+; CHECK-NEXT: v_mov_b32_e32 v243, 0
+; CHECK-NEXT: v_mov_b32_e32 v242, 0
+; CHECK-NEXT: v_mov_b32_e32 v241, 0
+; CHECK-NEXT: v_mov_b32_e32 v240, 0
+; CHECK-NEXT: v_mov_b32_e32 v247, 0
+; CHECK-NEXT: v_mov_b32_e32 v246, 0
+; CHECK-NEXT: v_mov_b32_e32 v245, 0
+; CHECK-NEXT: v_mov_b32_e32 v244, 0
+; CHECK-NEXT: v_mov_b32_e32 v251, 0
+; CHECK-NEXT: v_mov_b32_e32 v250, 0
+; CHECK-NEXT: v_mov_b32_e32 v249, 0
+; CHECK-NEXT: v_mov_b32_e32 v248, 0
+; CHECK-NEXT: v_mov_b32_e32 v255, 0
+; CHECK-NEXT: v_mov_b32_e32 v254, 0
+; CHECK-NEXT: v_mov_b32_e32 v253, 0
+; CHECK-NEXT: v_mov_b32_e32 v252, 0
+; CHECK-NEXT: v_mov_b32_e32 v239, 0
+; CHECK-NEXT: v_mov_b32_e32 v238, 0
+; CHECK-NEXT: v_mov_b32_e32 v237, 0
+; CHECK-NEXT: v_mov_b32_e32 v236, 0
+; CHECK-NEXT: v_mov_b32_e32 v235, 0
+; CHECK-NEXT: v_mov_b32_e32 v234, 0
+; CHECK-NEXT: v_mov_b32_e32 v233, 0
+; CHECK-NEXT: v_mov_b32_e32 v232, 0
+; CHECK-NEXT: v_mov_b32_e32 v231, 0
+; CHECK-NEXT: v_mov_b32_e32 v230, 0
+; CHECK-NEXT: v_mov_b32_e32 v229, 0
+; CHECK-NEXT: v_mov_b32_e32 v228, 0
+; CHECK-NEXT: v_mov_b32_e32 v227, 0
+; CHECK-NEXT: v_mov_b32_e32 v226, 0
+; CHECK-NEXT: v_mov_b32_e32 v225, 0
+; CHECK-NEXT: v_mov_b32_e32 v224, 0
+; CHECK-NEXT: v_accvgpr_write_b32 a0, 0
+; CHECK-NEXT: v_accvgpr_write_b32 a1, 0
+; CHECK-NEXT: v_accvgpr_write_b32 a2, 0
+; CHECK-NEXT: v_accvgpr_write_b32 a3, 0
+; CHECK-NEXT: s_waitcnt lgkmcnt(0)
+; CHECK-NEXT: v_accvgpr_write_b32 a4, s8
+; CHECK-NEXT: v_accvgpr_write_b32 a5, s9
+; CHECK-NEXT: v_accvgpr_write_b32 a6, s10
+; CHECK-NEXT: v_accvgpr_write_b32 a7, s11
+; CHECK-NEXT: v_accvgpr_write_b32 a8, s12
+; CHECK-NEXT: v_accvgpr_write_b32 a9, s13
+; CHECK-NEXT: v_accvgpr_write_b32 a10, s14
+; CHECK-NEXT: v_accvgpr_write_b32 a11, s15
+; CHECK-NEXT: .p2align 5, , 4
+; CHECK-NEXT: .LBB0_1: ; %loop
+; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
+; CHECK-NEXT: s_nop 1
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], a[4:7], a[8:11], a[0:3]
+; CHECK-NEXT: s_add_i32 s1, s1, 1
+; CHECK-NEXT: v_pk_add_f32 v[242:243], v[242:243], v[242:243]
+; CHECK-NEXT: v_pk_add_f32 v[240:241], v[240:241], v[240:241]
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], a[4:7], a[8:11], a[0:3]
+; CHECK-NEXT: v_add_f32_e64 v246, v246, v246
+; CHECK-NEXT: v_add_f32_e64 v247, v247, v247
+; CHECK-NEXT: v_pk_add_f32 v[244:245], v[244:245], v[244:245]
+; CHECK-NEXT: v_pk_add_f32 v[250:251], v[250:251], v[250:251]
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], a[4:7], a[8:11], a[0:3]
+; CHECK-NEXT: v_add_f32_e64 v248, v248, v248
+; CHECK-NEXT: v_add_f32_e64 v249, v249, v249
+; CHECK-NEXT: v_pk_add_f32 v[254:255], v[254:255], v[254:255]
+; CHECK-NEXT: v_pk_add_f32 v[252:253], v[252:253], v[252:253]
+; CHECK-NEXT: v_pk_add_f32 v[238:239], v[238:239], v[238:239]
+; CHECK-NEXT: v_pk_add_f32 v[236:237], v[236:237], v[236:237]
+; CHECK-NEXT: v_pk_add_f32 v[234:235], v[234:235], v[234:235]
+; CHECK-NEXT: v_pk_add_f32 v[232:233], v[232:233], v[232:233]
+; CHECK-NEXT: v_pk_add_f32 v[230:231], v[230:231], v[230:231]
+; CHECK-NEXT: v_pk_add_f32 v[228:229], v[228:229], v[228:229]
+; CHECK-NEXT: v_pk_add_f32 v[226:227], v[226:227], v[226:227]
+; CHECK-NEXT: v_pk_add_f32 v[224:225], v[224:225], v[224:225]
+; CHECK-NEXT: v_pk_add_f32 v[222:223], v[222:223], v[222:223]
+; CHECK-NEXT: v_pk_add_f32 v[220:221], v[220:221], v[220:221]
+; CHECK-NEXT: v_pk_add_f32 v[218:219], v[218:219], v[218:219]
+; CHECK-NEXT: v_pk_add_f32 v[216:217], v[216:217], v[216:217]
+; CHECK-NEXT: v_pk_add_f32 v[214:215], v[214:215], v[214:215]
+; CHECK-NEXT: v_pk_add_f32 v[212:213], v[212:213], v[212:213]
+; CHECK-NEXT: v_pk_add_f32 v[210:211], v[210:211], v[210:211]
+; CHECK-NEXT: v_pk_add_f32 v[208:209], v[208:209], v[208:209]
+; CHECK-NEXT: v_pk_add_f32 v[206:207], v[206:207], v[206:207]
+; CHECK-NEXT: v_pk_add_f32 v[204:205], v[204:205], v[204:205]
+; CHECK-NEXT: v_pk_add_f32 v[202:203], v[202:203], v[202:203]
+; CHECK-NEXT: v_pk_add_f32 v[200:201], v[200:201], v[200:201]
+; CHECK-NEXT: v_pk_add_f32 v[198:199], v[198:199], v[198:199]
+; CHECK-NEXT: v_pk_add_f32 v[196:197], v[196:197], v[196:197]
+; CHECK-NEXT: v_pk_add_f32 v[194:195], v[194:195], v[194:195]
+; CHECK-NEXT: v_pk_add_f32 v[192:193], v[192:193], v[192:193]
+; CHECK-NEXT: v_pk_add_f32 v[190:191], v[190:191], v[190:191]
+; CHECK-NEXT: v_pk_add_f32 v[188:189], v[188:189], v[188:189]
+; CHECK-NEXT: v_pk_add_f32 v[186:187], v[186:187], v[186:187]
+; CHECK-NEXT: v_pk_add_f32 v[184:185], v[184:185], v[184:185]
+; CHECK-NEXT: v_pk_add_f32 v[182:183], v[182:183], v[182:183]
+; CHECK-NEXT: v_pk_add_f32 v[180:181], v[180:181], v[180:181]
+; CHECK-NEXT: v_pk_add_f32 v[178:179], v[178:179], v[178:179]
+; CHECK-NEXT: v_pk_add_f32 v[176:177], v[176:177], v[176:177]
+; CHECK-NEXT: v_pk_add_f32 v[174:175], v[174:175], v[174:175]
+; CHECK-NEXT: v_pk_add_f32 v[172:173], v[172:173], v[172:173]
+; CHECK-NEXT: v_pk_add_f32 v[170:171], v[170:171], v[170:171]
+; CHECK-NEXT: v_pk_add_f32 v[168:169], v[168:169], v[168:169]
+; CHECK-NEXT: v_pk_add_f32 v[166:167], v[166:167], v[166:167]
+; CHECK-NEXT: v_pk_add_f32 v[164:165], v[164:165], v[164:165]
+; CHECK-NEXT: v_pk_add_f32 v[162:163], v[162:163], v[162:163]
+; CHECK-NEXT: v_pk_add_f32 v[160:161], v[160:161], v[160:161]
+; CHECK-NEXT: v_pk_add_f32 v[158:159], v[158:159], v[158:159]
+; CHECK-NEXT: v_pk_add_f32 v[156:157], v[156:157], v[156:157]
+; CHECK-NEXT: v_pk_add_f32 v[154:155], v[154:155], v[154:155]
+; CHECK-NEXT: v_pk_add_f32 v[152:153], v[152:153], v[152:153]
+; CHECK-NEXT: v_pk_add_f32 v[150:151], v[150:151], v[150:151]
+; CHECK-NEXT: v_pk_add_f32 v[148:149], v[148:149], v[148:149]
+; CHECK-NEXT: v_pk_add_f32 v[146:147], v[146:147], v[146:147]
+; CHECK-NEXT: v_pk_add_f32 v[144:145], v[144:145], v[144:145]
+; CHECK-NEXT: v_pk_add_f32 v[142:143], v[142:143], v[142:143]
+; CHECK-NEXT: v_pk_add_f32 v[140:141], v[140:141], v[140:141]
+; CHECK-NEXT: v_pk_add_f32 v[138:139], v[138:139], v[138:139]
+; CHECK-NEXT: v_pk_add_f32 v[136:137], v[136:137], v[136:137]
+; CHECK-NEXT: v_pk_add_f32 v[134:135], v[134:135], v[134:135]
+; CHECK-NEXT: v_pk_add_f32 v[132:133], v[132:133], v[132:133]
+; CHECK-NEXT: v_pk_add_f32 v[130:131], v[130:131], v[130:131]
+; CHECK-NEXT: v_pk_add_f32 v[128:129], v[128:129], v[128:129]
+; CHECK-NEXT: v_pk_add_f32 v[126:127], v[126:127], v[126:127]
+; CHECK-NEXT: v_pk_add_f32 v[124:125], v[124:125], v[124:125]
+; CHECK-NEXT: v_pk_add_f32 v[122:123], v[122:123], v[122:123]
+; CHECK-NEXT: v_pk_add_f32 v[120:121], v[120:121], v[120:121]
+; CHECK-NEXT: v_pk_add_f32 v[118:119], v[118:119], v[118:119]
+; CHECK-NEXT: v_pk_add_f32 v[116:117], v[116:117], v[116:117]
+; CHECK-NEXT: v_pk_add_f32 v[114:115], v[114:115], v[114:115]
+; CHECK-NEXT: v_pk_add_f32 v[112:113], v[112:113], v[112:113]
+; CHECK-NEXT: v_pk_add_f32 v[110:111], v[110:111], v[110:111]
+; CHECK-NEXT: v_pk_add_f32 v[108:109], v[108:109], v[108:109]
+; CHECK-NEXT: v_pk_add_f32 v[106:107], v[106:107], v[106:107]
+; CHECK-NEXT: v_pk_add_f32 v[104:105], v[104:105], v[104:105]
+; CHECK-NEXT: v_pk_add_f32 v[102:103], v[102:103], v[102:103]
+; CHECK-NEXT: v_pk_add_f32 v[100:101], v[100:101], v[100:101]
+; CHECK-NEXT: v_pk_add_f32 v[98:99], v[98:99], v[98:99]
+; CHECK-NEXT: v_pk_add_f32 v[96:97], v[96:97], v[96:97]
+; CHECK-NEXT: v_pk_add_f32 v[94:95], v[94:95], v[94:95]
+; CHECK-NEXT: v_pk_add_f32 v[92:93], v[92:93], v[92:93]
+; CHECK-NEXT: v_pk_add_f32 v[90:91], v[90:91], v[90:91]
+; CHECK-NEXT: v_pk_add_f32 v[88:89], v[88:89], v[88:89]
+; CHECK-NEXT: v_pk_add_f32 v[86:87], v[86:87], v[86:87]
+; CHECK-NEXT: v_pk_add_f32 v[84:85], v[84:85], v[84:85]
+; CHECK-NEXT: v_pk_add_f32 v[82:83], v[82:83], v[82:83]
+; CHECK-NEXT: v_pk_add_f32 v[80:81], v[80:81], v[80:81]
+; CHECK-NEXT: v_pk_add_f32 v[78:79], v[78:79], v[78:79]
+; CHECK-NEXT: v_pk_add_f32 v[76:77], v[76:77], v[76:77]
+; CHECK-NEXT: v_pk_add_f32 v[74:75], v[74:75], v[74:75]
+; CHECK-NEXT: v_pk_add_f32 v[72:73], v[72:73], v[72:73]
+; CHECK-NEXT: v_pk_add_f32 v[70:71], v[70:71], v[70:71]
+; CHECK-NEXT: v_pk_add_f32 v[68:69], v[68:69], v[68:69]
+; CHECK-NEXT: v_pk_add_f32 v[66:67], v[66:67], v[66:67]
+; CHECK-NEXT: v_pk_add_f32 v[64:65], v[64:65], v[64:65]
+; CHECK-NEXT: v_pk_add_f32 v[62:63], v[62:63], v[62:63]
+; CHECK-NEXT: v_pk_add_f32 v[60:61], v[60:61], v[60:61]
+; CHECK-NEXT: v_pk_add_f32 v[58:59], v[58:59], v[58:59]
+; CHECK-NEXT: v_pk_add_f32 v[56:57], v[56:57], v[56:57]
+; CHECK-NEXT: v_pk_add_f32 v[54:55], v[54:55], v[54:55]
+; CHECK-NEXT: v_pk_add_f32 v[52:53], v[52:53], v[52:53]
+; CHECK-NEXT: v_pk_add_f32 v[50:51], v[50:51], v[50:51]
+; CHECK-NEXT: v_pk_add_f32 v[48:49], v[48:49], v[48:49]
+; CHECK-NEXT: v_pk_add_f32 v[46:47], v[46:47], v[46:47]
+; CHECK-NEXT: v_pk_add_f32 v[44:45], v[44:45], v[44:45]
+; CHECK-NEXT: v_pk_add_f32 v[42:43], v[42:43], v[42:43]
+; CHECK-NEXT: v_pk_add_f32 v[40:41], v[40:41], v[40:41]
+; CHECK-NEXT: v_pk_add_f32 v[38:39], v[38:39], v[38:39]
+; CHECK-NEXT: v_pk_add_f32 v[36:37], v[36:37], v[36:37]
+; CHECK-NEXT: v_pk_add_f32 v[34:35], v[34:35], v[34:35]
+; CHECK-NEXT: v_pk_add_f32 v[32:33], v[32:33], v[32:33]
+; CHECK-NEXT: v_pk_add_f32 v[30:31], v[30:31], v[30:31]
+; CHECK-NEXT: v_pk_add_f32 v[28:29], v[28:29], v[28:29]
+; CHECK-NEXT: v_pk_add_f32 v[26:27], v[26:27], v[26:27]
+; CHECK-NEXT: v_pk_add_f32 v[24:25], v[24:25], v[24:25]
+; CHECK-NEXT: v_pk_add_f32 v[22:23], v[22:23], v[22:23]
+; CHECK-NEXT: v_pk_add_f32 v[20:21], v[20:21], v[20:21]
+; CHECK-NEXT: v_pk_add_f32 v[18:19], v[18:19], v[18:19]
+; CHECK-NEXT: v_pk_add_f32 v[16:17], v[16:17], v[16:17]
+; CHECK-NEXT: v_pk_add_f32 v[14:15], v[14:15], v[14:15]
+; CHECK-NEXT: v_pk_add_f32 v[12:13], v[12:13], v[12:13]
+; CHECK-NEXT: v_pk_add_f32 v[10:11], v[10:11], v[10:11]
+; CHECK-NEXT: v_pk_add_f32 v[8:9], v[8:9], v[8:9]
+; CHECK-NEXT: v_pk_add_f32 v[6:7], v[6:7], v[6:7]
+; CHECK-NEXT: v_pk_add_f32 v[4:5], v[4:5], v[4:5]
+; CHECK-NEXT: v_pk_add_f32 v[2:3], v[2:3], v[2:3]
+; CHECK-NEXT: s_cmp_lt_i32 s1, s0
+; CHECK-NEXT: v_pk_add_f32 v[0:1], v[0:1], v[0:1]
+; CHECK-NEXT: s_cbranch_scc1 .LBB0_1
+; CHECK-NEXT: ; %bb.2: ; %exit
+; CHECK-NEXT: v_accvgpr_write_b32 a4, s8
+; CHECK-NEXT: v_accvgpr_write_b32 a5, s9
+; CHECK-NEXT: v_accvgpr_write_b32 a6, s10
+; CHECK-NEXT: v_accvgpr_write_b32 a7, s11
+; CHECK-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x28
+; CHECK-NEXT: v_accvgpr_write_b32 a8, s12
+; CHECK-NEXT: v_accvgpr_write_b32 a9, s13
+; CHECK-NEXT: v_accvgpr_write_b32 a10, s14
+; CHECK-NEXT: v_accvgpr_write_b32 a11, s15
+; CHECK-NEXT: v_accvgpr_write_b32 a15, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a14, v2
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], a[4:7], a[8:11], a[0:3]
+; CHECK-NEXT: v_accvgpr_write_b32 a13, v1
+; CHECK-NEXT: v_accvgpr_write_b32 a12, v0
+; CHECK-NEXT: v_mov_b32_e32 v0, 0
+; CHECK-NEXT: s_waitcnt lgkmcnt(0)
+; CHECK-NEXT: global_store_dwordx4 v0, v[240:243], s[0:1] offset:2032
+; CHECK-NEXT: v_mov_b32_e32 v0, 0
+; CHECK-NEXT: global_store_dwordx4 v0, v[244:247], s[0:1] offset:2016
+; CHECK-NEXT: global_store_dwordx4 v0, v[248:251], s[0:1] offset:2000
+; CHECK-NEXT: v_mov_b32_e32 v0, 0
+; CHECK-NEXT: v_accvgpr_read_b32 v243, a3
+; CHECK-NEXT: v_accvgpr_read_b32 v242, a2
+; CHECK-NEXT: v_accvgpr_read_b32 v241, a1
+; CHECK-NEXT: v_accvgpr_read_b32 v240, a0
+; CHECK-NEXT: global_store_dwordx4 v0, v[252:255], s[0:1] offset:1984
+; CHECK-NEXT: s_nop 0
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[244:247], a[4:7], a[8:11], v[240:243]
+; CHECK-NEXT: s_nop 7
+; CHECK-NEXT: v_mov_b32_e32 v245, 0
+; CHECK-NEXT: global_store_dwordx4 v245, v[236:239], s[0:1] offset:1968
+; CHECK-NEXT: s_nop 1
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 v[236:239], a[8:11], a[4:7], v[240:243]
+; CHECK-NEXT: s_nop 7
+; CHECK-NEXT: v_add_f32_e32 v238, 1.0, v244
+; CHECK-NEXT: v_add_f32_e32 v239, 2.0, v236
+; CHECK-NEXT: global_store_dwordx2 v245, v[238:239], s[0:1]
+; CHECK-NEXT: global_store_dwordx4 v245, v[232:235], s[0:1] offset:1952
+; CHECK-NEXT: global_store_dwordx4 v245, v[228:231], s[0:1] offset:1936
+; CHECK-NEXT: global_store_dwordx4 v245, v[224:227], s[0:1] offset:1920
+; CHECK-NEXT: global_store_dwordx4 v245, v[220:223], s[0:1] offset:368
+; CHECK-NEXT: global_store_dwordx4 v245, v[216:219], s[0:1] offset:352
+; CHECK-NEXT: global_store_dwordx4 v245, v[212:215], s[0:1] offset:336
+; CHECK-NEXT: global_store_dwordx4 v245, v[208:211], s[0:1] offset:320
+; CHECK-NEXT: global_store_dwordx4 v245, v[204:207], s[0:1] offset:304
+; CHECK-NEXT: global_store_dwordx4 v245, v[200:203], s[0:1] offset:288
+; CHECK-NEXT: global_store_dwordx4 v245, v[196:199], s[0:1] offset:272
+; CHECK-NEXT: global_store_dwordx4 v245, v[192:195], s[0:1] offset:256
+; CHECK-NEXT: global_store_dwordx4 v245, v[188:191], s[0:1] offset:496
+; CHECK-NEXT: global_store_dwordx4 v245, v[184:187], s[0:1] offset:480
+; CHECK-NEXT: global_store_dwordx4 v245, v[180:183], s[0:1] offset:464
+; CHECK-NEXT: global_store_dwordx4 v245, v[176:179], s[0:1] offset:448
+; CHECK-NEXT: global_store_dwordx4 v245, v[172:175], s[0:1] offset:432
+; CHECK-NEXT: global_store_dwordx4 v245, v[168:171], s[0:1] offset:416
+; CHECK-NEXT: global_store_dwordx4 v245, v[164:167], s[0:1] offset:400
+; CHECK-NEXT: global_store_dwordx4 v245, v[160:163], s[0:1] offset:384
+; CHECK-NEXT: global_store_dwordx4 v245, v[156:159], s[0:1] offset:624
+; CHECK-NEXT: global_store_dwordx4 v245, v[152:155], s[0:1] offset:608
+; CHECK-NEXT: global_store_dwordx4 v245, v[148:151], s[0:1] offset:592
+; CHECK-NEXT: global_store_dwordx4 v245, v[144:147], s[0:1] offset:576
+; CHECK-NEXT: global_store_dwordx4 v245, v[140:143], s[0:1] offset:560
+; CHECK-NEXT: global_store_dwordx4 v245, v[136:139], s[0:1] offset:544
+; CHECK-NEXT: global_store_dwordx4 v245, v[132:135], s[0:1] offset:528
+; CHECK-NEXT: global_store_dwordx4 v245, v[128:131], s[0:1] offset:512
+; CHECK-NEXT: global_store_dwordx4 v245, v[124:127], s[0:1] offset:752
+; CHECK-NEXT: global_store_dwordx4 v245, v[120:123], s[0:1] offset:736
+; CHECK-NEXT: global_store_dwordx4 v245, v[116:119], s[0:1] offset:720
+; CHECK-NEXT: global_store_dwordx4 v245, v[112:115], s[0:1] offset:704
+; CHECK-NEXT: global_store_dwordx4 v245, v[108:111], s[0:1] offset:688
+; CHECK-NEXT: global_store_dwordx4 v245, v[104:107], s[0:1] offset:672
+; CHECK-NEXT: global_store_dwordx4 v245, v[100:103], s[0:1] offset:656
+; CHECK-NEXT: global_store_dwordx4 v245, v[96:99], s[0:1] offset:640
+; CHECK-NEXT: global_store_dwordx4 v245, v[92:95], s[0:1] offset:880
+; CHECK-NEXT: global_store_dwordx4 v245, v[88:91], s[0:1] offset:864
+; CHECK-NEXT: global_store_dwordx4 v245, v[84:87], s[0:1] offset:848
+; CHECK-NEXT: global_store_dwordx4 v245, v[80:83], s[0:1] offset:832
+; CHECK-NEXT: global_store_dwordx4 v245, v[76:79], s[0:1] offset:816
+; CHECK-NEXT: global_store_dwordx4 v245, v[72:75], s[0:1] offset:800
+; CHECK-NEXT: global_store_dwordx4 v245, v[68:71], s[0:1] offset:784
+; CHECK-NEXT: global_store_dwordx4 v245, v[64:67], s[0:1] offset:768
+; CHECK-NEXT: global_store_dwordx4 v245, v[60:63], s[0:1] offset:1008
+; CHECK-NEXT: global_store_dwordx4 v245, v[56:59], s[0:1] offset:992
+; CHECK-NEXT: global_store_dwordx4 v245, v[52:55], s[0:1] offset:976
+; CHECK-NEXT: global_store_dwordx4 v245, v[48:51], s[0:1] offset:960
+; CHECK-NEXT: global_store_dwordx4 v245, v[44:47], s[0:1] offset:944
+; CHECK-NEXT: global_store_dwordx4 v245, v[40:43], s[0:1] offset:928
+; CHECK-NEXT: global_store_dwordx4 v245, v[36:39], s[0:1] offset:912
+; CHECK-NEXT: global_store_dwordx4 v245, v[32:35], s[0:1] offset:896
+; CHECK-NEXT: global_store_dwordx4 v245, v[28:31], s[0:1] offset:1136
+; CHECK-NEXT: global_store_dwordx4 v245, v[24:27], s[0:1] offset:1120
+; CHECK-NEXT: global_store_dwordx4 v245, v[20:23], s[0:1] offset:1104
+; CHECK-NEXT: global_store_dwordx4 v245, v[16:19], s[0:1] offset:1088
+; CHECK-NEXT: global_store_dwordx4 v245, v[12:15], s[0:1] offset:1072
+; CHECK-NEXT: global_store_dwordx4 v245, v[8:11], s[0:1] offset:1056
+; CHECK-NEXT: global_store_dwordx4 v245, v[4:7], s[0:1] offset:1040
+; CHECK-NEXT: global_store_dwordx4 v245, a[12:15], s[0:1] offset:1024
+; CHECK-NEXT: s_endpgm
+; Loop body: MFMAs in AGPR form.
+; Exit block: one MFMA consuming the loop result (AGPR form), then
+; exactly 4 v_accvgpr_read (one 128-bit copy), shared by both vgprcd MFMAs.
+
+ <8 x half> %a, <8 x half> %b, i32 %n, ptr addrspace(1) %out
+) #0 {
+entry:
+ br label %loop
+
+loop:
+ %iv = phi i32 [ 0, %entry ], [ %iv.next, %loop ]
+ %acc = phi <4 x float> [ zeroinitializer, %entry ], [ %m2, %loop ]
+ ; 8 loop-carried <32 x float> vectors = 256 VGPRs of pressure.
+ %vc0 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc0n, %loop ]
+ %vc1 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc1n, %loop ]
+ %vc2 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc2n, %loop ]
+ %vc3 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc3n, %loop ]
+ %vc4 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc4n, %loop ]
+ %vc5 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc5n, %loop ]
+ %vc6 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc6n, %loop ]
+ %vc7 = phi <32 x float> [ zeroinitializer, %entry ], [ %vc7n, %loop ]
+
+ %m0 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a, <8 x half> %b, <4 x float> %acc, i32 0, i32 0, i32 0)
+ %m1 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a, <8 x half> %b, <4 x float> %m0, i32 0, i32 0, i32 0)
+ %m2 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a, <8 x half> %b, <4 x float> %m1, i32 0, i32 0, i32 0)
+
+ ; Keep carriers alive.
+ %vc0n = fadd <32 x float> %vc0, %vc0
+ %vc1n = fadd <32 x float> %vc1, %vc1
+ %vc2n = fadd <32 x float> %vc2, %vc2
+ %vc3n = fadd <32 x float> %vc3, %vc3
+ %vc4n = fadd <32 x float> %vc4, %vc4
+ %vc5n = fadd <32 x float> %vc5, %vc5
+ %vc6n = fadd <32 x float> %vc6, %vc6
+ %vc7n = fadd <32 x float> %vc7, %vc7
+
+ %iv.next = add i32 %iv, 1
+ %cmp = icmp slt i32 %iv.next, %n
+ br i1 %cmp, label %loop, label %exit
+
+exit:
+ %m3 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a, <8 x half> %b, <4 x float> %m2, i32 0, i32 0, i32 0)
+ %out1 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a, <8 x half> %b, <4 x float> %m3, i32 0, i32 0, i32 0)
+ %out2 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %b, <8 x half> %a, <4 x float> %m3, i32 0, i32 0, i32 0)
+
+ ; VGPR-requiring uses of the exit MFMAs to force accvgpr_read copies.
+ %e0 = extractelement <4 x float> %out1, i32 0
+ %e1 = extractelement <4 x float> %out2, i32 0
+ %add0 = fadd float %e0, 1.0
+ %add1 = fadd float %e1, 2.0
+
+ ; Store results + carriers to keep everything live.
+ store float %add0, ptr addrspace(1) %out, align 4
+ %gep1 = getelementptr float, ptr addrspace(1) %out, i32 1
+ store float %add1, ptr addrspace(1) %gep1, align 4
+ %gep2 = getelementptr <32 x float>, ptr addrspace(1) %out, i32 15
+ store <32 x float> %vc0n, ptr addrspace(1) %gep2, align 128
+ %gep3 = getelementptr <32 x float>, ptr addrspace(1) %out, i32 2
+ store <32 x float> %vc1n, ptr addrspace(1) %gep3, align 128
+ %gep4 = getelementptr <32 x float>, ptr addrspace(1) %out, i32 3
+ store <32 x float> %vc2n, ptr addrspace(1) %gep4, align 128
+ %gep5 = getelementptr <32 x float>, ptr addrspace(1) %out, i32 4
+ store <32 x float> %vc3n, ptr addrspace(1) %gep5, align 128
+ %gep6 = getelementptr <32 x float>, ptr addrspace(1) %out, i32 5
+ store <32 x float> %vc4n, ptr addrspace(1) %gep6, align 128
+ %gep7 = getelementptr <32 x float>, ptr addrspace(1) %out, i32 6
+ store <32 x float> %vc5n, ptr addrspace(1) %gep7, align 128
+ %gep8 = getelementptr <32 x float>, ptr addrspace(1) %out, i32 7
+ store <32 x float> %vc6n, ptr addrspace(1) %gep8, align 128
+ %gep9 = getelementptr <32 x float>, ptr addrspace(1) %out, i32 8
+ store <32 x float> %vc7n, ptr addrspace(1) %gep9, align 128
+ ret void
+}
+
+attributes #0 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" }
More information about the llvm-commits
mailing list