[llvm] [AMDGPU] Add RewriteMFMAForm scheduling stage tests (PR #196008)
via llvm-commits
llvm-commits at lists.llvm.org
Tue May 5 23:47:23 PDT 2026
https://github.com/xgxanq created https://github.com/llvm/llvm-project/pull/196008
Add lit tests for the RewriteMFMAFormStage machine scheduler pass:
- rewrite-mfma-form-ReachingUseTracker.ll: cross-block use via ReachingUseTracker (Case 2/3)
- rewrite-mfma-form-case2-case3.ll: Case 2 (AGPR->VGPR at use) and Case 3 (VGPR->AGPR at def)
- rewrite-mfma-form-check-half-rewrite.mir: partial rewrite cost gating
- rewrite-mfma-form-v7slice-pattern.ll: v7_slice pattern on gfx950 mfma.f32.16x16x32.f16
>From 5d8ee787bdc4c5392baf737539c99f6a10bc1e19 Mon Sep 17 00:00:00 2001
From: anqfu <anqfu at amd.com>
Date: Wed, 6 May 2026 06:32:45 +0000
Subject: [PATCH] [AMDGPU] Add RewriteMFMAForm scheduling stage tests
Add lit tests for the RewriteMFMAFormStage machine scheduler pass:
- rewrite-mfma-form-ReachingUseTracker.ll: cross-block use via ReachingUseTracker (Case 2/3)
- rewrite-mfma-form-case2-case3.ll: Case 2 (AGPR->VGPR at use) and Case 3 (VGPR->AGPR at def)
- rewrite-mfma-form-check-half-rewrite.mir: partial rewrite cost gating
- rewrite-mfma-form-v7slice-pattern.ll: v7_slice pattern on gfx950 mfma.f32.16x16x32.f16
---
.../rewrite-mfma-form-ReachingUseTracker.ll | 512 ++++++++++++++++++
.../AMDGPU/rewrite-mfma-form-case2-case3.ll | 143 +++++
.../rewrite-mfma-form-check-half-rewrite.mir | 84 +++
.../rewrite-mfma-form-v7slice-pattern.ll | 332 ++++++++++++
4 files changed, 1071 insertions(+)
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-ReachingUseTracker.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-case2-case3.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-check-half-rewrite.mir
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-v7slice-pattern.ll
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-ReachingUseTracker.ll b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-ReachingUseTracker.ll
new file mode 100644
index 0000000000000..7ee7627424861
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-ReachingUseTracker.ll
@@ -0,0 +1,512 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a \
+; RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+; RUN: < %s | FileCheck %s
+;
+; Verify RewriteMFMAFormStage output:
+; Case 3: v_accvgpr_write inserted in entry block — zeroinitializer acc
+; converted VGPR->AGPR before first loop iteration.
+; Loop: all 9 MFMAs rewritten to AGPR form (a[...] C/D operands).
+; Case 2: v_accvgpr_read inserted at exit block entry (ReachingUseTracker)
+; to convert AGPR MFMA results -> VGPR before fptrunc (V_CVT_F16_F32).
+;
+; CHECK-LABEL: test_ReachingUseTracker_crossblock_use:
+; CHECK: ; @test_ReachingUseTracker_crossblock_use
+; Case 3: v_accvgpr_write_b32 inserted in entry — zeroinitializer reaching-defs
+; of MFMA src2 converted VGPR->AGPR before first loop iteration.
+; CHECK: ; %bb.0: ; %entry
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; All 9 MFMAs rewritten to AGPR form (a[...] C/D). Loop-carried acc stays in AGPR.
+; CHECK: .LBB0_1: ; %loop.body
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; Case 2: v_accvgpr_read_b32 inserted at exit block entry (ReachingUseTracker) —
+; AGPR MFMA results converted AGPR->VGPR before fptrunc (V_CVT_F16_F32).
+; CHECK: ; %bb.2: ; %exit
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+
+define amdgpu_kernel void @test_ReachingUseTracker_crossblock_use(
+ ptr addrspace(1) %out,
+ <2 x i16> %a,
+ <2 x i16> %b,
+ i32 %n) #0 {
+entry:
+ br label %loop.body
+
+loop.body:
+ %i = phi i32 [ 0, %entry ], [ %i.next, %loop.body ]
+ ; 8 loop-carried <32 x float>: 8*32=256 VGPRs non-MFMA pressure.
+ ; Back-edge: self-insert (no dep on MFMA results, never DCE'd due to store use).
+ %c0 = phi <32 x float> [ zeroinitializer, %entry ], [ %c0n, %loop.body ]
+ %c1 = phi <32 x float> [ zeroinitializer, %entry ], [ %c1n, %loop.body ]
+ %c2 = phi <32 x float> [ zeroinitializer, %entry ], [ %c2n, %loop.body ]
+ %c3 = phi <32 x float> [ zeroinitializer, %entry ], [ %c3n, %loop.body ]
+ %c4 = phi <32 x float> [ zeroinitializer, %entry ], [ %c4n, %loop.body ]
+ %c5 = phi <32 x float> [ zeroinitializer, %entry ], [ %c5n, %loop.body ]
+ %c6 = phi <32 x float> [ zeroinitializer, %entry ], [ %c6n, %loop.body ]
+ %c7 = phi <32 x float> [ zeroinitializer, %entry ], [ %c7n, %loop.body ]
+ ; 9 MFMA accumulators: 9*4=36 VGPRs.
+ %acc0 = phi <4 x float> [ zeroinitializer, %entry ], [ %r0, %loop.body ]
+ %acc1 = phi <4 x float> [ zeroinitializer, %entry ], [ %r1, %loop.body ]
+ %acc2 = phi <4 x float> [ zeroinitializer, %entry ], [ %r2, %loop.body ]
+ %acc3 = phi <4 x float> [ zeroinitializer, %entry ], [ %r3, %loop.body ]
+ %acc4 = phi <4 x float> [ zeroinitializer, %entry ], [ %r4, %loop.body ]
+ %acc5 = phi <4 x float> [ zeroinitializer, %entry ], [ %r5, %loop.body ]
+ %acc6 = phi <4 x float> [ zeroinitializer, %entry ], [ %r6, %loop.body ]
+ %acc7 = phi <4 x float> [ zeroinitializer, %entry ], [ %r7, %loop.body ]
+ %acc8 = phi <4 x float> [ zeroinitializer, %entry ], [ %r8, %loop.body ]
+
+ ; 9 MFMAs. Results only used in exit (CopyForUse exit only, UseFreq~0).
+ %r0 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc0, i32 0, i32 0, i32 0)
+ %r1 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc1, i32 0, i32 0, i32 0)
+ %r2 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc2, i32 0, i32 0, i32 0)
+ %r3 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc3, i32 0, i32 0, i32 0)
+ %r4 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc4, i32 0, i32 0, i32 0)
+ %r5 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc5, i32 0, i32 0, i32 0)
+ %r6 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc6, i32 0, i32 0, i32 0)
+ %r7 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc7, i32 0, i32 0, i32 0)
+ %r8 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc8, i32 0, i32 0, i32 0)
+
+ ; Variable index prevents folding of carrier self-inserts.
+ %eidx = and i32 %i, 31
+ ; Carrier self-insert: no dep on MFMA results. GPU scheduler fills MFMA
+ ; latency with these, naturally scheduling them AFTER MFMAs -> carriers live
+ ; across all MFMAs -> peak pressure 256+36=292 VGPRs.
+ %c0e = extractelement <32 x float> %c0, i32 0
+ %c1e = extractelement <32 x float> %c1, i32 0
+ %c2e = extractelement <32 x float> %c2, i32 0
+ %c3e = extractelement <32 x float> %c3, i32 0
+ %c4e = extractelement <32 x float> %c4, i32 0
+ %c5e = extractelement <32 x float> %c5, i32 0
+ %c6e = extractelement <32 x float> %c6, i32 0
+ %c7e = extractelement <32 x float> %c7, i32 0
+ %c0n = insertelement <32 x float> %c0, float %c0e, i32 %eidx
+ %c1n = insertelement <32 x float> %c1, float %c1e, i32 %eidx
+ %c2n = insertelement <32 x float> %c2, float %c2e, i32 %eidx
+ %c3n = insertelement <32 x float> %c3, float %c3e, i32 %eidx
+ %c4n = insertelement <32 x float> %c4, float %c4e, i32 %eidx
+ %c5n = insertelement <32 x float> %c5, float %c5e, i32 %eidx
+ %c6n = insertelement <32 x float> %c6, float %c6e, i32 %eidx
+ %c7n = insertelement <32 x float> %c7, float %c7e, i32 %eidx
+
+ ; Store one carrier element (keeps carriers live, prevents DCE).
+ %csum = fadd float %c0e, %c1e
+ store float %csum, ptr addrspace(1) %out, align 4
+
+ %i.next = add i32 %i, 1
+ %cond = icmp eq i32 %i.next, %n
+ br i1 %cond, label %exit, label %loop.body
+
+exit:
+ ; Cross-block non-MAI uses: MFMAs in loop.body, fptrunc uses in exit.
+ ; -> ReachingUseTracker.
+ %e0_0 = extractelement <4 x float> %r0, i32 0
+ %e1_0 = extractelement <4 x float> %r0, i32 1
+ %v0a = insertelement <2 x float> poison, float %e0_0, i32 0
+ %v0 = insertelement <2 x float> %v0a, float %e1_0, i32 1
+ %h0 = fptrunc <2 x float> %v0 to <2 x half>
+ %p_h0 = getelementptr i16, ptr addrspace(1) %out, i32 2
+ store <2 x half> %h0, ptr addrspace(1) %p_h0, align 2
+
+ %e0_1 = extractelement <4 x float> %r1, i32 0
+ %e1_1 = extractelement <4 x float> %r1, i32 1
+ %v1a = insertelement <2 x float> poison, float %e0_1, i32 0
+ %v1 = insertelement <2 x float> %v1a, float %e1_1, i32 1
+ %h1 = fptrunc <2 x float> %v1 to <2 x half>
+ %p_h1 = getelementptr i16, ptr addrspace(1) %out, i32 4
+ store <2 x half> %h1, ptr addrspace(1) %p_h1, align 2
+
+ %e0_2 = extractelement <4 x float> %r2, i32 0
+ %e1_2 = extractelement <4 x float> %r2, i32 1
+ %v2a = insertelement <2 x float> poison, float %e0_2, i32 0
+ %v2 = insertelement <2 x float> %v2a, float %e1_2, i32 1
+ %h2 = fptrunc <2 x float> %v2 to <2 x half>
+ %p_h2 = getelementptr i16, ptr addrspace(1) %out, i32 6
+ store <2 x half> %h2, ptr addrspace(1) %p_h2, align 2
+
+ %e0_3 = extractelement <4 x float> %r3, i32 0
+ %e1_3 = extractelement <4 x float> %r3, i32 1
+ %v3a = insertelement <2 x float> poison, float %e0_3, i32 0
+ %v3 = insertelement <2 x float> %v3a, float %e1_3, i32 1
+ %h3 = fptrunc <2 x float> %v3 to <2 x half>
+ %p_h3 = getelementptr i16, ptr addrspace(1) %out, i32 8
+ store <2 x half> %h3, ptr addrspace(1) %p_h3, align 2
+
+ %e0_4 = extractelement <4 x float> %r4, i32 0
+ %e1_4 = extractelement <4 x float> %r4, i32 1
+ %v4a = insertelement <2 x float> poison, float %e0_4, i32 0
+ %v4 = insertelement <2 x float> %v4a, float %e1_4, i32 1
+ %h4 = fptrunc <2 x float> %v4 to <2 x half>
+ %p_h4 = getelementptr i16, ptr addrspace(1) %out, i32 10
+ store <2 x half> %h4, ptr addrspace(1) %p_h4, align 2
+
+ %e0_5 = extractelement <4 x float> %r5, i32 0
+ %e1_5 = extractelement <4 x float> %r5, i32 1
+ %v5a = insertelement <2 x float> poison, float %e0_5, i32 0
+ %v5 = insertelement <2 x float> %v5a, float %e1_5, i32 1
+ %h5 = fptrunc <2 x float> %v5 to <2 x half>
+ %p_h5 = getelementptr i16, ptr addrspace(1) %out, i32 12
+ store <2 x half> %h5, ptr addrspace(1) %p_h5, align 2
+
+ %e0_6 = extractelement <4 x float> %r6, i32 0
+ %e1_6 = extractelement <4 x float> %r6, i32 1
+ %v6a = insertelement <2 x float> poison, float %e0_6, i32 0
+ %v6 = insertelement <2 x float> %v6a, float %e1_6, i32 1
+ %h6 = fptrunc <2 x float> %v6 to <2 x half>
+ %p_h6 = getelementptr i16, ptr addrspace(1) %out, i32 14
+ store <2 x half> %h6, ptr addrspace(1) %p_h6, align 2
+
+ %e0_7 = extractelement <4 x float> %r7, i32 0
+ %e1_7 = extractelement <4 x float> %r7, i32 1
+ %v7a = insertelement <2 x float> poison, float %e0_7, i32 0
+ %v7 = insertelement <2 x float> %v7a, float %e1_7, i32 1
+ %h7 = fptrunc <2 x float> %v7 to <2 x half>
+ %p_h7 = getelementptr i16, ptr addrspace(1) %out, i32 16
+ store <2 x half> %h7, ptr addrspace(1) %p_h7, align 2
+
+ %e0_8 = extractelement <4 x float> %r8, i32 0
+ %e1_8 = extractelement <4 x float> %r8, i32 1
+ %v8a = insertelement <2 x float> poison, float %e0_8, i32 0
+ %v8 = insertelement <2 x float> %v8a, float %e1_8, i32 1
+ %h8 = fptrunc <2 x float> %v8 to <2 x half>
+ %p_h8 = getelementptr i16, ptr addrspace(1) %out, i32 18
+ store <2 x half> %h8, ptr addrspace(1) %p_h8, align 2
+
+ ret void
+}
+
+attributes #0 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" }
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-case2-case3.ll b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-case2-case3.ll
new file mode 100644
index 0000000000000..e7f7b0d77e053
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-case2-case3.ll
@@ -0,0 +1,143 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a \
+; RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+; RUN: < %s | FileCheck %s
+
+define amdgpu_kernel void @test_case2_case3(
+; CHECK-LABEL: test_case2_case3:
+; CHECK: ; @test_case2_case3
+; Case 3: v_accvgpr_write in entry. Loop MFMAs all-AGPR. Case 2: v_accvgpr_read in exit.
+; CHECK: ; %bb.0: ; %entry
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, 0
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, 0
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, 0
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, 0
+; CHECK: v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: .LBB0_1: ; %loop.body
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_4x4x2bf16 a[{{[0-9:]+}}], a{{[0-9]+}}, a{{[0-9]+}}, a[{{[0-9:]+}}]
+; CHECK: ; %bb.2: ; %exit
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+ ptr addrspace(1) %out,
+ <2 x i16> %a,
+ <2 x i16> %b,
+ i32 %n) #0 {
+entry:
+ br label %loop.body
+
+loop.body:
+ %i = phi i32 [ 0, %entry ], [ %i.next, %loop.body ]
+ ; 8 loop-carried <32 x float>: 8*32=256 VGPRs non-MFMA pressure.
+ %c0 = phi <32 x float> [ zeroinitializer, %entry ], [ %c0n, %loop.body ]
+ %c1 = phi <32 x float> [ zeroinitializer, %entry ], [ %c1n, %loop.body ]
+ %c2 = phi <32 x float> [ zeroinitializer, %entry ], [ %c2n, %loop.body ]
+ %c3 = phi <32 x float> [ zeroinitializer, %entry ], [ %c3n, %loop.body ]
+ %c4 = phi <32 x float> [ zeroinitializer, %entry ], [ %c4n, %loop.body ]
+ %c5 = phi <32 x float> [ zeroinitializer, %entry ], [ %c5n, %loop.body ]
+ %c6 = phi <32 x float> [ zeroinitializer, %entry ], [ %c6n, %loop.body ]
+ %c7 = phi <32 x float> [ zeroinitializer, %entry ], [ %c7n, %loop.body ]
+ ; 9 MFMA accumulators: 9*4=36 VGPRs. Case 3: entry zeroinitializer -> AGPR.
+ %acc0 = phi <4 x float> [ zeroinitializer, %entry ], [ %r0, %loop.body ]
+ %acc1 = phi <4 x float> [ zeroinitializer, %entry ], [ %r1, %loop.body ]
+ %acc2 = phi <4 x float> [ zeroinitializer, %entry ], [ %r2, %loop.body ]
+ %acc3 = phi <4 x float> [ zeroinitializer, %entry ], [ %r3, %loop.body ]
+ %acc4 = phi <4 x float> [ zeroinitializer, %entry ], [ %r4, %loop.body ]
+ %acc5 = phi <4 x float> [ zeroinitializer, %entry ], [ %r5, %loop.body ]
+ %acc6 = phi <4 x float> [ zeroinitializer, %entry ], [ %r6, %loop.body ]
+ %acc7 = phi <4 x float> [ zeroinitializer, %entry ], [ %r7, %loop.body ]
+ %acc8 = phi <4 x float> [ zeroinitializer, %entry ], [ %r8, %loop.body ]
+
+ ; 9 MFMAs — results only used in exit (Case 2: cross-block non-MAI use).
+ %r0 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc0, i32 0, i32 0, i32 0)
+ %r1 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc1, i32 0, i32 0, i32 0)
+ %r2 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc2, i32 0, i32 0, i32 0)
+ %r3 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc3, i32 0, i32 0, i32 0)
+ %r4 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc4, i32 0, i32 0, i32 0)
+ %r5 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc5, i32 0, i32 0, i32 0)
+ %r6 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc6, i32 0, i32 0, i32 0)
+ %r7 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc7, i32 0, i32 0, i32 0)
+ %r8 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %acc8, i32 0, i32 0, i32 0)
+
+ ; Variable index prevents folding: keeps carriers live across all MFMAs.
+ %eidx = and i32 %i, 31
+ %c0e = extractelement <32 x float> %c0, i32 0
+ %c1e = extractelement <32 x float> %c1, i32 0
+ %c2e = extractelement <32 x float> %c2, i32 0
+ %c3e = extractelement <32 x float> %c3, i32 0
+ %c4e = extractelement <32 x float> %c4, i32 0
+ %c5e = extractelement <32 x float> %c5, i32 0
+ %c6e = extractelement <32 x float> %c6, i32 0
+ %c7e = extractelement <32 x float> %c7, i32 0
+ %c0n = insertelement <32 x float> %c0, float %c0e, i32 %eidx
+ %c1n = insertelement <32 x float> %c1, float %c1e, i32 %eidx
+ %c2n = insertelement <32 x float> %c2, float %c2e, i32 %eidx
+ %c3n = insertelement <32 x float> %c3, float %c3e, i32 %eidx
+ %c4n = insertelement <32 x float> %c4, float %c4e, i32 %eidx
+ %c5n = insertelement <32 x float> %c5, float %c5e, i32 %eidx
+ %c6n = insertelement <32 x float> %c6, float %c6e, i32 %eidx
+ %c7n = insertelement <32 x float> %c7, float %c7e, i32 %eidx
+
+ %csum = fadd float %c0e, %c1e
+ store float %csum, ptr addrspace(1) %out, align 4
+
+ %i.next = add i32 %i, 1
+ %cond = icmp eq i32 %i.next, %n
+ br i1 %cond, label %exit, label %loop.body
+
+exit:
+ ; Case 2: non-MAI uses of MFMA dst in exit block (different block from loop.body).
+ ; Case 3: reaching def of each %acc in entry is zeroinitializer (non-MAI).
+ ; Expected: v_accvgpr_read (AGPR->VGPR) inserted before these extractelements.
+ %e0 = extractelement <4 x float> %r0, i32 0
+ %e1 = extractelement <4 x float> %r1, i32 0
+ store float %e0, ptr addrspace(1) %out, align 4
+ %p1 = getelementptr float, ptr addrspace(1) %out, i32 1
+ store float %e1, ptr addrspace(1) %p1, align 4
+ ret void
+}
+
+attributes #0 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" }
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-check-half-rewrite.mir b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-check-half-rewrite.mir
new file mode 100644
index 0000000000000..55d080dd272f4
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-check-half-rewrite.mir
@@ -0,0 +1,84 @@
+# XFAIL: *
+# REQUIRES: asserts
+# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 \
+# RUN: -run-pass=machine-scheduler \
+# RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+# RUN: -filetype=null %s
+#
+# Test: Bug 2 in RewriteMFMAFormStage — physical register in MFMA src2.
+#
+# Root cause (GCNSchedStrategy.cpp):
+# rewrite() processes RewriteCands in iteration order. For each MFMA:
+# MI->setDesc(ReplacementOp); // opcode mutated
+# if (!Src2Reg.isVirtual()) return false; // Bug 2: mid-loop bail-out
+# A return false here leaves all previously-processed MFMAs in AGPR form
+# while subsequent MFMAs are partially mutated — inconsistent MIR state.
+#
+# Trigger:
+# MFMA_B uses $vgpr4_vgpr5_vgpr6_vgpr7 (physical VGPR) as src2.
+# Without the fix: initHeuristics calls findReachingDefs($vgpr4) ->
+# LIS->getInterval($vgpr4) -> virtRegIndex() asserts "Not a virtual register".
+#
+# Without the fix: llc crashes (SIGABRT) -> non-zero exit -> RUN fails -> XFAIL.
+# With the fix: llc exits 0 -> RUN passes -> XPASS (remove XFAIL, convert test).
+
+--- |
+ define void @test_bug2_physreg_src2() #0 {
+ entry:
+ unreachable
+ }
+
+ attributes #0 = { "amdgpu-waves-per-eu"="1,1" "amdgpu-flat-work-group-size"="64,64" }
+...
+
+---
+name: test_bug2_physreg_src2
+tracksRegLiveness: true
+liveins:
+ - { reg: '$vgpr4' }
+ - { reg: '$vgpr5' }
+ - { reg: '$vgpr6' }
+ - { reg: '$vgpr7' }
+machineFunctionInfo:
+ isEntryFunction: true
+ scratchRSrcReg: '$sgpr96_sgpr97_sgpr98_sgpr99'
+ stackPtrOffsetReg: '$sgpr32'
+ argumentInfo:
+ privateSegmentBuffer: { reg: '$sgpr0_sgpr1_sgpr2_sgpr3' }
+ kernargSegmentPtr: { reg: '$sgpr4_sgpr5' }
+ workGroupIDX: { reg: '$sgpr6' }
+ privateSegmentWaveByteOffset: { reg: '$sgpr7' }
+ workItemIDX: { reg: '$vgpr0' }
+ sgprForEXECCopy: '$sgpr100_sgpr101'
+body: |
+ bb.0:
+ liveins: $vgpr0, $sgpr4_sgpr5, $vgpr4, $vgpr5, $vgpr6, $vgpr7
+
+ ; 9 x vreg_1024 = 288 VGPRs > 256: triggers RegionsWithExcessArchVGPR.
+ %0:vreg_1024 = IMPLICIT_DEF
+ %1:vreg_1024 = IMPLICIT_DEF
+ %2:vreg_1024 = IMPLICIT_DEF
+ %3:vreg_1024 = IMPLICIT_DEF
+ %4:vreg_1024 = IMPLICIT_DEF
+ %5:vreg_1024 = IMPLICIT_DEF
+ %6:vreg_1024 = IMPLICIT_DEF
+ %7:vreg_1024 = IMPLICIT_DEF
+ %8:vreg_1024 = IMPLICIT_DEF
+
+ %10:av_128_align2 = IMPLICIT_DEF
+ %11:av_128_align2 = IMPLICIT_DEF
+ %12:vreg_64_align2 = IMPLICIT_DEF
+ %13:vgpr_32 = IMPLICIT_DEF
+
+ SCHED_BARRIER 0
+
+ ; MFMA_A: all virtual registers.
+ %acc_a:vreg_128_align2 = IMPLICIT_DEF
+ %res_a:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 %10:av_128_align2, %11:av_128_align2, %acc_a:vreg_128_align2, 4, 4, %12.sub0:vreg_64_align2, %13:vgpr_32, 0, 0, implicit $mode, implicit $exec
+
+ ; MFMA_B: src2 = $vgpr4_vgpr5_vgpr6_vgpr7 (physical VGPR) — triggers crash.
+ %res_b:vreg_128_align2 = contract nofpexcept V_MFMA_SCALE_F32_16X16X128_F8F6F4_f4_f4_vgprcd_e64 %10:av_128_align2, %11:av_128_align2, $vgpr4_vgpr5_vgpr6_vgpr7, 4, 4, %12.sub0:vreg_64_align2, %13:vgpr_32, 0, 0, implicit $mode, implicit $exec
+
+ KILL %0, %1, %2, %3, %4, %5, %6, %7, %8, %res_a, %res_b
+
+ S_ENDPGM 0
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-v7slice-pattern.ll b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-v7slice-pattern.ll
new file mode 100644
index 0000000000000..02a5360f1acce
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-mfma-form-v7slice-pattern.ll
@@ -0,0 +1,332 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 \
+; RUN: -amdgpu-disable-rewrite-mfma-form-sched-stage=false \
+; RUN: < %s | FileCheck %s
+;
+; Test: RewriteMFMAFormStage — v7_slice pattern (gfx950, mfma.f32.16x16x32.f16)
+;
+; Distilled from v7_slice.llir (Triton matmul kernel, gfx950).
+;
+; Key structural features preserved from v7_slice:
+; (1) Loop body (%loop) accumulates via loop-carried SCALAR float phis,
+; NOT <4 x float> phis. Each MFMA acc is built by insertelement from
+; 4 individual float phis then used as src2/dst in MFMAs within the loop.
+; (2) Epilogue (%epilogue) ALSO contains MFMAs: the final K-tile iteration
+; is peeled. Epilogue MFMAs take scalar float phis from %loop as acc,
+; then produce results consumed by fptrunc+store — Case 2.
+; (3) Scalar float phis initialized to 0.0 in entry (non-MAI def) — Case 3.
+;
+; CFG: entry -> loop (back-edge) -> epilogue -> ret
+;
+; Pressure design (gfx950, 256 ArchVGPR limit):
+; 16 loop-carried scalar float phis (4 chains x 4 elems) = 16 VGPRs.
+; 4 MFMA chains of 2, each acc = 4 float phis = 16 VGPRs in loop.
+; 4 x <8 x half> operands (a0, a1, b0, b1) = 16 VGPRs.
+; 8 x <32 x float> loop-carried vec carriers (vc0-vc7) = 256 VGPRs.
+; Peak ArchVGPR = 256 + 16 + 16 + 16 = 304 > 256 -> RegionsWithExcessArchVGPR.
+; After rewrite: scalar phis=16 VGPR, <8 x half>=16 VGPR, vec=256 VGPR, MFMA acc->AGPR.
+; getRewriteCost() < 0 -> rewrite() fires.
+;
+; Expected:
+; Case 3: v_accvgpr_write inserted after insertelement defs of MFMA acc
+; (entry zeroinitializer -> scalar float phi -> insertelement = non-MAI)
+; Case 2: v_accvgpr_read before fptrunc in epilogue block
+; MFMA opcodes: v_mfma_f32_16x16x32_f16 vgprcd -> agprcd form
+
+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 @test_v7slice_scalar_phi_acc(
+; CHECK-LABEL: test_v7slice_scalar_phi_acc:
+; CHECK: ; @test_v7slice_scalar_phi_acc
+; Case 3: v_accvgpr_write inserted for scalar-phi acc (non-MAI def -> AGPR init).
+; Loop body MFMAs use AGPR C/D (a[...], a[...], a[...], a[...]).
+; CHECK: ; %bb.0: ; %entry
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
+; CHECK: .LBB0_1: ; %loop
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; Case 2: epilogue MFMAs take VGPR src0/src1, AGPR C/D. v_accvgpr_read before fptrunc.
+; CHECK: ; %bb.2: ; %epilogue
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_mfma_f32_16x16x32_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+; CHECK: v_accvgpr_read_b32 v{{[0-9]+}}, a{{[0-9]+}}
+ ptr addrspace(1) %out,
+ <8 x half> %a0,
+ <8 x half> %a1,
+ <8 x half> %b0,
+ <8 x half> %b1,
+ i32 %n) #0 {
+entry:
+ br label %loop
+
+loop:
+ %i = phi i32 [ 0, %entry ], [ %i.next, %loop ]
+
+ ; -------------------------------------------------------------------
+ ; 8 x <32 x float> loop-carried vector carriers = 256 VGPRs.
+ ; Self-insert with variable index prevents folding.
+ ; Kept live by store use in loop body.
+ ; -------------------------------------------------------------------
+ %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 ]
+
+ ; -------------------------------------------------------------------
+ ; 16 loop-carried scalar float phis = 16 VGPRs.
+ ; These are the per-element accumulator slots, mirroring v7_slice's
+ ; pattern where each output element is a separate loop-carried float.
+ ; Initialized to 0.0 in entry (non-MAI def) -> Case 3 triggers on the
+ ; insertelement that builds the MFMA src2 from these scalars.
+ ; -------------------------------------------------------------------
+ ; MFMA chain A acc (4 floats):
+ %a0_0 = phi float [ 0.0, %entry ], [ %ra0_0, %loop ]
+ %a0_1 = phi float [ 0.0, %entry ], [ %ra0_1, %loop ]
+ %a0_2 = phi float [ 0.0, %entry ], [ %ra0_2, %loop ]
+ %a0_3 = phi float [ 0.0, %entry ], [ %ra0_3, %loop ]
+ ; MFMA chain B acc (4 floats):
+ %a1_0 = phi float [ 0.0, %entry ], [ %ra1_0, %loop ]
+ %a1_1 = phi float [ 0.0, %entry ], [ %ra1_1, %loop ]
+ %a1_2 = phi float [ 0.0, %entry ], [ %ra1_2, %loop ]
+ %a1_3 = phi float [ 0.0, %entry ], [ %ra1_3, %loop ]
+ ; MFMA chain C acc (4 floats):
+ %a2_0 = phi float [ 0.0, %entry ], [ %ra2_0, %loop ]
+ %a2_1 = phi float [ 0.0, %entry ], [ %ra2_1, %loop ]
+ %a2_2 = phi float [ 0.0, %entry ], [ %ra2_2, %loop ]
+ %a2_3 = phi float [ 0.0, %entry ], [ %ra2_3, %loop ]
+ ; MFMA chain D acc (4 floats):
+ %a3_0 = phi float [ 0.0, %entry ], [ %ra3_0, %loop ]
+ %a3_1 = phi float [ 0.0, %entry ], [ %ra3_1, %loop ]
+ %a3_2 = phi float [ 0.0, %entry ], [ %ra3_2, %loop ]
+ %a3_3 = phi float [ 0.0, %entry ], [ %ra3_3, %loop ]
+
+ ; -------------------------------------------------------------------
+ ; Build <4 x float> acc vectors from scalar phi elements (v7_slice pattern).
+ ; Each insertelement is a non-MAI def of an element of the acc register.
+ ; -------------------------------------------------------------------
+ %accA0 = insertelement <4 x float> poison, float %a0_0, i32 0
+ %accA1 = insertelement <4 x float> %accA0, float %a0_1, i32 1
+ %accA2 = insertelement <4 x float> %accA1, float %a0_2, i32 2
+ %accA = insertelement <4 x float> %accA2, float %a0_3, i32 3
+
+ %accB0 = insertelement <4 x float> poison, float %a1_0, i32 0
+ %accB1 = insertelement <4 x float> %accB0, float %a1_1, i32 1
+ %accB2 = insertelement <4 x float> %accB1, float %a1_2, i32 2
+ %accB = insertelement <4 x float> %accB2, float %a1_3, i32 3
+
+ %accC0 = insertelement <4 x float> poison, float %a2_0, i32 0
+ %accC1 = insertelement <4 x float> %accC0, float %a2_1, i32 1
+ %accC2 = insertelement <4 x float> %accC1, float %a2_2, i32 2
+ %accC = insertelement <4 x float> %accC2, float %a2_3, i32 3
+
+ %accD0 = insertelement <4 x float> poison, float %a3_0, i32 0
+ %accD1 = insertelement <4 x float> %accD0, float %a3_1, i32 1
+ %accD2 = insertelement <4 x float> %accD1, float %a3_2, i32 2
+ %accD = insertelement <4 x float> %accD2, float %a3_3, i32 3
+
+ ; -------------------------------------------------------------------
+ ; MFMA chains in loop body (2 MFMAs per chain, chained acc).
+ ; Results are loop-carried back via extractelement -> scalar phi.
+ ; -------------------------------------------------------------------
+ %rA_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %accA, i32 0, i32 0, i32 0)
+ %rA = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %rA_v, i32 0, i32 0, i32 0)
+
+ %rB_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %accB, i32 0, i32 0, i32 0)
+ %rB = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %rB_v, i32 0, i32 0, i32 0)
+
+ %rC_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %accC, i32 0, i32 0, i32 0)
+ %rC = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %rC_v, i32 0, i32 0, i32 0)
+
+ %rD_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %accD, i32 0, i32 0, i32 0)
+ %rD = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %rD_v, i32 0, i32 0, i32 0)
+
+ ; Extract scalar results to carry back via phi.
+ %ra0_0 = extractelement <4 x float> %rA, i32 0
+ %ra0_1 = extractelement <4 x float> %rA, i32 1
+ %ra0_2 = extractelement <4 x float> %rA, i32 2
+ %ra0_3 = extractelement <4 x float> %rA, i32 3
+
+ %ra1_0 = extractelement <4 x float> %rB, i32 0
+ %ra1_1 = extractelement <4 x float> %rB, i32 1
+ %ra1_2 = extractelement <4 x float> %rB, i32 2
+ %ra1_3 = extractelement <4 x float> %rB, i32 3
+
+ %ra2_0 = extractelement <4 x float> %rC, i32 0
+ %ra2_1 = extractelement <4 x float> %rC, i32 1
+ %ra2_2 = extractelement <4 x float> %rC, i32 2
+ %ra2_3 = extractelement <4 x float> %rC, i32 3
+
+ %ra3_0 = extractelement <4 x float> %rD, i32 0
+ %ra3_1 = extractelement <4 x float> %rD, i32 1
+ %ra3_2 = extractelement <4 x float> %rD, i32 2
+ %ra3_3 = extractelement <4 x float> %rD, i32 3
+
+ ; Vector carrier self-inserts (prevent register reuse).
+ %eidx = and i32 %i, 31
+ %vc0e = extractelement <32 x float> %vc0, i32 0
+ %vc1e = extractelement <32 x float> %vc1, i32 0
+ %vc2e = extractelement <32 x float> %vc2, i32 0
+ %vc3e = extractelement <32 x float> %vc3, i32 0
+ %vc4e = extractelement <32 x float> %vc4, i32 0
+ %vc5e = extractelement <32 x float> %vc5, i32 0
+ %vc6e = extractelement <32 x float> %vc6, i32 0
+ %vc7e = extractelement <32 x float> %vc7, i32 0
+ %vc0n = insertelement <32 x float> %vc0, float %vc0e, i32 %eidx
+ %vc1n = insertelement <32 x float> %vc1, float %vc1e, i32 %eidx
+ %vc2n = insertelement <32 x float> %vc2, float %vc2e, i32 %eidx
+ %vc3n = insertelement <32 x float> %vc3, float %vc3e, i32 %eidx
+ %vc4n = insertelement <32 x float> %vc4, float %vc4e, i32 %eidx
+ %vc5n = insertelement <32 x float> %vc5, float %vc5e, i32 %eidx
+ %vc6n = insertelement <32 x float> %vc6, float %vc6e, i32 %eidx
+ %vc7n = insertelement <32 x float> %vc7, float %vc7e, i32 %eidx
+
+ %vcs = fadd float %vc0e, %vc1e
+ store float %vcs, ptr addrspace(1) %out, align 4
+
+ %i.next = add i32 %i, 1
+ %cond = icmp eq i32 %i.next, %n
+ br i1 %cond, label %epilogue, label %loop
+
+epilogue:
+ ; -------------------------------------------------------------------
+ ; v7_slice pattern: epilogue ALSO has MFMAs (peeled final K-tile).
+ ; Acc built from the same scalar float phis coming out of %loop.
+ ; These insertelements are non-MAI defs -> Case 3.
+ ; -------------------------------------------------------------------
+ %eaccA0 = insertelement <4 x float> poison, float %ra0_0, i32 0
+ %eaccA1 = insertelement <4 x float> %eaccA0, float %ra0_1, i32 1
+ %eaccA2 = insertelement <4 x float> %eaccA1, float %ra0_2, i32 2
+ %eaccA = insertelement <4 x float> %eaccA2, float %ra0_3, i32 3
+
+ %eaccB0 = insertelement <4 x float> poison, float %ra1_0, i32 0
+ %eaccB1 = insertelement <4 x float> %eaccB0, float %ra1_1, i32 1
+ %eaccB2 = insertelement <4 x float> %eaccB1, float %ra1_2, i32 2
+ %eaccB = insertelement <4 x float> %eaccB2, float %ra1_3, i32 3
+
+ %eaccC0 = insertelement <4 x float> poison, float %ra2_0, i32 0
+ %eaccC1 = insertelement <4 x float> %eaccC0, float %ra2_1, i32 1
+ %eaccC2 = insertelement <4 x float> %eaccC1, float %ra2_2, i32 2
+ %eaccC = insertelement <4 x float> %eaccC2, float %ra2_3, i32 3
+
+ %eaccD0 = insertelement <4 x float> poison, float %ra3_0, i32 0
+ %eaccD1 = insertelement <4 x float> %eaccD0, float %ra3_1, i32 1
+ %eaccD2 = insertelement <4 x float> %eaccD1, float %ra3_2, i32 2
+ %eaccD = insertelement <4 x float> %eaccD2, float %ra3_3, i32 3
+
+ ; Epilogue MFMAs (2 per chain, as in v7_slice).
+ %erA_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %eaccA, i32 0, i32 0, i32 0)
+ %erA = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %erA_v, i32 0, i32 0, i32 0)
+
+ %erB_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %eaccB, i32 0, i32 0, i32 0)
+ %erB = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %erB_v, i32 0, i32 0, i32 0)
+
+ %erC_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %eaccC, i32 0, i32 0, i32 0)
+ %erC = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %erC_v, i32 0, i32 0, i32 0)
+
+ %erD_v = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a0, <8 x half> %b0, <4 x float> %eaccD, i32 0, i32 0, i32 0)
+ %erD = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %a1, <8 x half> %b1, <4 x float> %erD_v, i32 0, i32 0, i32 0)
+
+ ; -------------------------------------------------------------------
+ ; Non-MAI use of epilogue MFMA results via fptrunc (Case 2).
+ ; v7_slice: results shuffled and fptrunc'd to fp16 for stores.
+ ; -------------------------------------------------------------------
+ %shA = shufflevector <4 x float> %erA, <4 x float> poison, <2 x i32> <i32 0, i32 1>
+ %hA = fptrunc <2 x float> %shA to <2 x half>
+ store <2 x half> %hA, ptr addrspace(1) %out, align 2
+
+ %shB = shufflevector <4 x float> %erB, <4 x float> poison, <2 x i32> <i32 0, i32 1>
+ %hB = fptrunc <2 x float> %shB to <2 x half>
+ %pB = getelementptr i16, ptr addrspace(1) %out, i32 2
+ store <2 x half> %hB, ptr addrspace(1) %pB, align 2
+
+ %shC = shufflevector <4 x float> %erC, <4 x float> poison, <2 x i32> <i32 0, i32 1>
+ %hC = fptrunc <2 x float> %shC to <2 x half>
+ %pC = getelementptr i16, ptr addrspace(1) %out, i32 4
+ store <2 x half> %hC, ptr addrspace(1) %pC, align 2
+
+ %shD = shufflevector <4 x float> %erD, <4 x float> poison, <2 x i32> <i32 0, i32 1>
+ %hD = fptrunc <2 x float> %shD to <2 x half>
+ %pD = getelementptr i16, ptr addrspace(1) %out, i32 6
+ store <2 x half> %hD, ptr addrspace(1) %pD, align 2
+
+ 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