[llvm] 4c31d2c - AMDGPU: Make rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll less allocator sensitive (#226913)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Sep 28 03:32:54 PDT 2026
Author: Matt Arsenault
Date: 2026-09-28T12:32:47+02:00
New Revision: 4c31d2c7641e180c2c9098e466b2247bb470fc25
URL: https://github.com/llvm/llvm-project/commit/4c31d2c7641e180c2c9098e466b2247bb470fc25
DIFF: https://github.com/llvm/llvm-project/commit/4c31d2c7641e180c2c9098e466b2247bb470fc25.diff
LOG: AMDGPU: Make rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll less allocator sensitive (#226913)
This test is sensitive to the exact split and spills which occur, and disappeared under
a future upstream improvement. Use basic RA with a fixed occupancy since it more
stably produces the spill pattern.
Also add a codegen reference test for the same kernel, so future codegen
improvements are visible.
Co-authored-by: Claude Opus 5 <noreply at anthropic.com>
Added:
llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store-codegen.ll
Modified:
llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
Removed:
################################################################################
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store-codegen.ll b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store-codegen.ll
new file mode 100644
index 0000000000000..987b827a48664
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store-codegen.ll
@@ -0,0 +1,694 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -mtriple=amdgpu9.0a-amd-amdhsa -O3 -amdgpu-use-amdgpu-trackers=1 < %s | FileCheck %s
+
+; Codegen reference for the kernel in
+; rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll, which is sensitive to whether
+; the register allocator ends up spilling.
+
+define amdgpu_kernel void @multi_store_spill_slot() #0 {
+; CHECK-LABEL: multi_store_spill_slot:
+; CHECK: ; %bb.0: ; %entry
+; CHECK-NEXT: s_mov_b32 s0, 0
+; CHECK-NEXT: v_mov_b32_e32 v3, 0
+; CHECK-NEXT: s_mov_b32 s1, s0
+; CHECK-NEXT: v_mov_b32_e32 v28, 0x3c003c00
+; CHECK-NEXT: v_pk_mov_b32 v[30:31], s[0:1], s[0:1] op_sel:[0,1]
+; CHECK-NEXT: v_mov_b32_e32 v0, 0x7fc00000
+; CHECK-NEXT: v_mov_b32_e32 v213, 1.0
+; CHECK-NEXT: v_mov_b32_e32 v29, v28
+; CHECK-NEXT: s_and_b64 vcc, exec, -1
+; CHECK-NEXT: v_mov_b32_e32 v2, v3
+; CHECK-NEXT: v_mov_b32_e32 v70, v3
+; CHECK-NEXT: v_mov_b32_e32 v20, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a64, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a1, v3
+; CHECK-NEXT: v_mov_b32_e32 v86, v3
+; CHECK-NEXT: v_mov_b32_e32 v162, v3
+; CHECK-NEXT: v_mov_b32_e32 v50, v3
+; CHECK-NEXT: v_mov_b32_e32 v212, v3
+; CHECK-NEXT: v_mov_b32_e32 v67, v3
+; CHECK-NEXT: v_mov_b32_e32 v66, v3
+; CHECK-NEXT: v_mov_b32_e32 v196, v3
+; CHECK-NEXT: v_mov_b32_e32 v146, v3
+; CHECK-NEXT: v_mov_b32_e32 v32, v3
+; CHECK-NEXT: v_mov_b32_e32 v243, v3
+; CHECK-NEXT: v_mov_b32_e32 v114, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a17, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a16, v3
+; CHECK-NEXT: v_mov_b32_e32 v145, v3
+; CHECK-NEXT: v_mov_b32_e32 v34, v3
+; CHECK-NEXT: v_mov_b32_e32 v35, v3
+; CHECK-NEXT: .LBB0_1: ; %do.body
+; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
+; CHECK-NEXT: v_mov_b32_e32 v4, v3
+; CHECK-NEXT: v_mov_b32_e32 v5, v3
+; CHECK-NEXT: v_mov_b32_e32 v6, v3
+; CHECK-NEXT: v_mov_b32_e32 v7, v3
+; CHECK-NEXT: v_mov_b32_e32 v8, v3
+; CHECK-NEXT: v_mov_b32_e32 v9, v3
+; CHECK-NEXT: v_mov_b32_e32 v10, v3
+; CHECK-NEXT: v_mov_b32_e32 v11, v3
+; CHECK-NEXT: v_mov_b32_e32 v12, v3
+; CHECK-NEXT: v_mov_b32_e32 v13, v3
+; CHECK-NEXT: v_mov_b32_e32 v14, v3
+; CHECK-NEXT: v_mov_b32_e32 v15, v3
+; CHECK-NEXT: v_mov_b32_e32 v16, v3
+; CHECK-NEXT: v_mov_b32_e32 v17, v3
+; CHECK-NEXT: v_mov_b32_e32 v214, v213
+; CHECK-NEXT: v_mov_b32_e32 v215, v213
+; CHECK-NEXT: v_mov_b32_e32 v216, v213
+; CHECK-NEXT: v_mov_b32_e32 v217, v213
+; CHECK-NEXT: v_mov_b32_e32 v218, v213
+; CHECK-NEXT: v_mov_b32_e32 v219, v213
+; CHECK-NEXT: v_mov_b32_e32 v220, v213
+; CHECK-NEXT: v_mov_b32_e32 v221, v213
+; CHECK-NEXT: v_mov_b32_e32 v222, v213
+; CHECK-NEXT: v_mov_b32_e32 v223, v213
+; CHECK-NEXT: v_mov_b32_e32 v224, v213
+; CHECK-NEXT: v_mov_b32_e32 v225, v213
+; CHECK-NEXT: v_mov_b32_e32 v226, v213
+; CHECK-NEXT: v_mov_b32_e32 v227, v213
+; CHECK-NEXT: v_accvgpr_write_b32 a47, v17
+; CHECK-NEXT: v_accvgpr_write_b32 a46, v16
+; CHECK-NEXT: v_accvgpr_write_b32 a45, v15
+; CHECK-NEXT: v_accvgpr_write_b32 a44, v14
+; CHECK-NEXT: v_accvgpr_write_b32 a43, v13
+; CHECK-NEXT: v_accvgpr_write_b32 a42, v12
+; CHECK-NEXT: v_accvgpr_write_b32 a41, v11
+; CHECK-NEXT: v_accvgpr_write_b32 a40, v10
+; CHECK-NEXT: v_accvgpr_write_b32 a39, v9
+; CHECK-NEXT: v_accvgpr_write_b32 a38, v8
+; CHECK-NEXT: v_accvgpr_write_b32 a37, v7
+; CHECK-NEXT: v_accvgpr_write_b32 a36, v6
+; CHECK-NEXT: v_accvgpr_write_b32 a35, v5
+; CHECK-NEXT: v_accvgpr_write_b32 a34, v4
+; CHECK-NEXT: v_accvgpr_write_b32 a33, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a32, v2
+; CHECK-NEXT: v_accvgpr_write_b32 a48, v212
+; CHECK-NEXT: v_accvgpr_write_b32 a49, v213
+; CHECK-NEXT: v_accvgpr_write_b32 a50, v214
+; CHECK-NEXT: v_accvgpr_write_b32 a51, v215
+; CHECK-NEXT: v_accvgpr_write_b32 a52, v216
+; CHECK-NEXT: v_accvgpr_write_b32 a53, v217
+; CHECK-NEXT: v_accvgpr_write_b32 a54, v218
+; CHECK-NEXT: v_accvgpr_write_b32 a55, v219
+; CHECK-NEXT: v_accvgpr_write_b32 a56, v220
+; CHECK-NEXT: v_accvgpr_write_b32 a57, v221
+; CHECK-NEXT: v_accvgpr_write_b32 a58, v222
+; CHECK-NEXT: v_accvgpr_write_b32 a59, v223
+; CHECK-NEXT: v_accvgpr_write_b32 a60, v224
+; CHECK-NEXT: v_accvgpr_write_b32 a61, v225
+; CHECK-NEXT: v_accvgpr_write_b32 a62, v226
+; CHECK-NEXT: v_accvgpr_write_b32 a63, v227
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[32:47], v[30:31], v[30:31], a[32:47]
+; CHECK-NEXT: v_mov_b32_e32 v13, v70
+; CHECK-NEXT: v_mov_b32_e32 v68, v3
+; CHECK-NEXT: v_mov_b32_e32 v69, v3
+; CHECK-NEXT: v_mov_b32_e32 v70, v3
+; CHECK-NEXT: v_mov_b32_e32 v71, v3
+; CHECK-NEXT: v_mov_b32_e32 v72, v3
+; CHECK-NEXT: v_mov_b32_e32 v73, v3
+; CHECK-NEXT: v_mov_b32_e32 v74, v3
+; CHECK-NEXT: v_mov_b32_e32 v75, v3
+; CHECK-NEXT: v_mov_b32_e32 v76, v3
+; CHECK-NEXT: v_mov_b32_e32 v77, v3
+; CHECK-NEXT: v_mov_b32_e32 v78, v3
+; CHECK-NEXT: v_mov_b32_e32 v79, v3
+; CHECK-NEXT: v_mov_b32_e32 v80, v3
+; CHECK-NEXT: v_mov_b32_e32 v81, v3
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[48:63], v[30:31], v[30:31], a[48:63]
+; CHECK-NEXT: v_mov_b32_e32 v242, v32
+; CHECK-NEXT: v_mov_b32_e32 v1, v3
+; CHECK-NEXT: v_mov_b32_e32 v2, v3
+; CHECK-NEXT: v_mov_b32_e32 v14, v20
+; CHECK-NEXT: v_mov_b32_e32 v228, v3
+; CHECK-NEXT: v_mov_b32_e32 v229, v3
+; CHECK-NEXT: v_mov_b32_e32 v230, v3
+; CHECK-NEXT: v_mov_b32_e32 v231, v3
+; CHECK-NEXT: v_mov_b32_e32 v232, v3
+; CHECK-NEXT: v_mov_b32_e32 v233, v3
+; CHECK-NEXT: v_mov_b32_e32 v234, v3
+; CHECK-NEXT: v_mov_b32_e32 v235, v3
+; CHECK-NEXT: v_mov_b32_e32 v236, v3
+; CHECK-NEXT: v_mov_b32_e32 v237, v3
+; CHECK-NEXT: v_mov_b32_e32 v238, v3
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[66:81], v[30:31], v[30:31], v[66:81]
+; CHECK-NEXT: v_mov_b32_e32 v239, v3
+; CHECK-NEXT: v_mov_b32_e32 v240, v3
+; CHECK-NEXT: v_mov_b32_e32 v241, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a18, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a19, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a20, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a21, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a22, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a23, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a24, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a25, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a26, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a27, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a28, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a29, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a30, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a31, v3
+; CHECK-NEXT: v_accvgpr_mov_b32 a0, a64
+; CHECK-NEXT: v_accvgpr_write_b32 a95, v15
+; CHECK-NEXT: v_accvgpr_write_b32 a94, v14
+; CHECK-NEXT: v_accvgpr_write_b32 a93, v13
+; CHECK-NEXT: v_accvgpr_write_b32 a92, v12
+; CHECK-NEXT: v_accvgpr_write_b32 a91, v11
+; CHECK-NEXT: v_accvgpr_write_b32 a90, v10
+; CHECK-NEXT: v_accvgpr_write_b32 a89, v9
+; CHECK-NEXT: v_accvgpr_write_b32 a88, v8
+; CHECK-NEXT: v_accvgpr_write_b32 a87, v7
+; CHECK-NEXT: v_accvgpr_write_b32 a86, v6
+; CHECK-NEXT: v_accvgpr_write_b32 a85, v5
+; CHECK-NEXT: v_accvgpr_write_b32 a84, v4
+; CHECK-NEXT: v_accvgpr_write_b32 a83, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a82, v2
+; CHECK-NEXT: v_accvgpr_write_b32 a81, v1
+; CHECK-NEXT: v_accvgpr_write_b32 a80, v0
+; CHECK-NEXT: v_mov_b32_e32 v5, v86
+; CHECK-NEXT: v_mov_b32_e32 v13, v3
+; CHECK-NEXT: v_mov_b32_e32 v14, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a64, v228
+; CHECK-NEXT: v_accvgpr_write_b32 a65, v229
+; CHECK-NEXT: v_accvgpr_write_b32 a66, v230
+; CHECK-NEXT: v_accvgpr_write_b32 a67, v231
+; CHECK-NEXT: v_accvgpr_write_b32 a68, v232
+; CHECK-NEXT: v_accvgpr_write_b32 a69, v233
+; CHECK-NEXT: v_accvgpr_write_b32 a70, v234
+; CHECK-NEXT: v_accvgpr_write_b32 a71, v235
+; CHECK-NEXT: v_accvgpr_write_b32 a72, v236
+; CHECK-NEXT: v_accvgpr_write_b32 a73, v237
+; CHECK-NEXT: v_accvgpr_write_b32 a74, v238
+; CHECK-NEXT: v_accvgpr_write_b32 a75, v239
+; CHECK-NEXT: v_accvgpr_write_b32 a76, v240
+; CHECK-NEXT: v_accvgpr_write_b32 a77, v241
+; CHECK-NEXT: v_accvgpr_write_b32 a78, v242
+; CHECK-NEXT: v_accvgpr_write_b32 a79, v243
+; CHECK-NEXT: v_accvgpr_mov_b32 a111, a47
+; CHECK-NEXT: v_accvgpr_mov_b32 a110, a46
+; CHECK-NEXT: v_accvgpr_mov_b32 a109, a45
+; CHECK-NEXT: v_accvgpr_mov_b32 a108, a44
+; CHECK-NEXT: v_accvgpr_mov_b32 a107, a43
+; CHECK-NEXT: v_accvgpr_mov_b32 a106, a42
+; CHECK-NEXT: v_accvgpr_mov_b32 a105, a41
+; CHECK-NEXT: v_accvgpr_mov_b32 a104, a40
+; CHECK-NEXT: v_accvgpr_mov_b32 a103, a39
+; CHECK-NEXT: v_accvgpr_mov_b32 a102, a38
+; CHECK-NEXT: v_accvgpr_mov_b32 a101, a37
+; CHECK-NEXT: v_accvgpr_mov_b32 a100, a36
+; CHECK-NEXT: v_accvgpr_mov_b32 a99, a35
+; CHECK-NEXT: v_accvgpr_mov_b32 a98, a34
+; CHECK-NEXT: v_accvgpr_mov_b32 a97, a33
+; CHECK-NEXT: v_accvgpr_mov_b32 a96, a32
+; CHECK-NEXT: v_accvgpr_mov_b32 a32, a48
+; CHECK-NEXT: v_accvgpr_mov_b32 a33, a49
+; CHECK-NEXT: v_accvgpr_mov_b32 a34, a50
+; CHECK-NEXT: v_accvgpr_mov_b32 a35, a51
+; CHECK-NEXT: v_accvgpr_mov_b32 a36, a52
+; CHECK-NEXT: v_accvgpr_mov_b32 a37, a53
+; CHECK-NEXT: v_accvgpr_mov_b32 a38, a54
+; CHECK-NEXT: v_accvgpr_mov_b32 a39, a55
+; CHECK-NEXT: v_accvgpr_mov_b32 a40, a56
+; CHECK-NEXT: v_accvgpr_mov_b32 a41, a57
+; CHECK-NEXT: v_accvgpr_mov_b32 a42, a58
+; CHECK-NEXT: v_accvgpr_mov_b32 a43, a59
+; CHECK-NEXT: v_accvgpr_mov_b32 a44, a60
+; CHECK-NEXT: v_accvgpr_mov_b32 a45, a61
+; CHECK-NEXT: v_accvgpr_mov_b32 a46, a62
+; CHECK-NEXT: v_accvgpr_mov_b32 a47, a63
+; CHECK-NEXT: v_accvgpr_mov_b32 a63, a31
+; CHECK-NEXT: v_accvgpr_mov_b32 a62, a30
+; CHECK-NEXT: v_accvgpr_mov_b32 a61, a29
+; CHECK-NEXT: v_accvgpr_mov_b32 a60, a28
+; CHECK-NEXT: v_accvgpr_mov_b32 a59, a27
+; CHECK-NEXT: v_accvgpr_mov_b32 a58, a26
+; CHECK-NEXT: v_accvgpr_mov_b32 a57, a25
+; CHECK-NEXT: v_accvgpr_mov_b32 a56, a24
+; CHECK-NEXT: v_accvgpr_mov_b32 a55, a23
+; CHECK-NEXT: v_accvgpr_mov_b32 a54, a22
+; CHECK-NEXT: v_accvgpr_mov_b32 a53, a21
+; CHECK-NEXT: v_accvgpr_mov_b32 a52, a20
+; CHECK-NEXT: v_accvgpr_mov_b32 a51, a19
+; CHECK-NEXT: v_accvgpr_mov_b32 a50, a18
+; CHECK-NEXT: v_accvgpr_mov_b32 a49, a17
+; CHECK-NEXT: v_accvgpr_mov_b32 a48, a16
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[64:79], v[30:31], v[30:31], a[64:79]
+; CHECK-NEXT: v_pk_mov_b32 v[100:101], v[16:17], v[16:17] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[98:99], v[14:15], v[14:15] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[96:97], v[12:13], v[12:13] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[94:95], v[10:11], v[10:11] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[92:93], v[8:9], v[8:9] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[90:91], v[6:7], v[6:7] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[88:89], v[4:5], v[4:5] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[86:87], v[2:3], v[2:3] op_sel:[0,1]
+; CHECK-NEXT: v_mov_b32_e32 v36, v3
+; CHECK-NEXT: v_mov_b32_e32 v37, v3
+; CHECK-NEXT: v_mov_b32_e32 v38, v3
+; CHECK-NEXT: v_mov_b32_e32 v39, v3
+; CHECK-NEXT: v_mov_b32_e32 v40, v3
+; CHECK-NEXT: v_mov_b32_e32 v41, v3
+; CHECK-NEXT: v_mov_b32_e32 v42, v3
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[48:63], v[30:31], v[30:31], a[48:63]
+; CHECK-NEXT: v_mov_b32_e32 v43, v3
+; CHECK-NEXT: v_mov_b32_e32 v44, v3
+; CHECK-NEXT: v_mov_b32_e32 v45, v3
+; CHECK-NEXT: v_mov_b32_e32 v46, v3
+; CHECK-NEXT: v_mov_b32_e32 v47, v3
+; CHECK-NEXT: v_mov_b32_e32 v48, v3
+; CHECK-NEXT: v_mov_b32_e32 v49, v3
+; CHECK-NEXT: v_pk_mov_b32 v[228:229], v[48:49], v[48:49] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[226:227], v[46:47], v[46:47] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[224:225], v[44:45], v[44:45] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[222:223], v[42:43], v[42:43] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[220:221], v[40:41], v[40:41] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[218:219], v[38:39], v[38:39] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[216:217], v[36:37], v[36:37] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[214:215], v[34:35], v[34:35] op_sel:[0,1]
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[4:19], v[30:31], v[30:31], v[66:81]
+; CHECK-NEXT: v_pk_mov_b32 v[32:33], v[214:215], v[214:215] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[34:35], v[216:217], v[216:217] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[36:37], v[218:219], v[218:219] op_sel:[0,1]
+; CHECK-NEXT: v_accvgpr_mov_b32 a127, a63
+; CHECK-NEXT: v_accvgpr_mov_b32 a126, a62
+; CHECK-NEXT: v_accvgpr_mov_b32 a125, a61
+; CHECK-NEXT: v_accvgpr_mov_b32 a124, a60
+; CHECK-NEXT: v_accvgpr_mov_b32 a123, a59
+; CHECK-NEXT: v_accvgpr_mov_b32 a122, a58
+; CHECK-NEXT: v_accvgpr_mov_b32 a121, a57
+; CHECK-NEXT: v_accvgpr_mov_b32 a120, a56
+; CHECK-NEXT: v_accvgpr_mov_b32 a119, a55
+; CHECK-NEXT: v_accvgpr_mov_b32 a118, a54
+; CHECK-NEXT: v_accvgpr_mov_b32 a117, a53
+; CHECK-NEXT: v_accvgpr_mov_b32 a116, a52
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[80:95], v[30:31], v[30:31], a[80:95]
+; CHECK-NEXT: v_accvgpr_mov_b32 a115, a51
+; CHECK-NEXT: v_accvgpr_mov_b32 a114, a50
+; CHECK-NEXT: v_accvgpr_mov_b32 a113, a49
+; CHECK-NEXT: v_accvgpr_mov_b32 a112, a48
+; CHECK-NEXT: v_accvgpr_mov_b32 a48, a64
+; CHECK-NEXT: v_accvgpr_mov_b32 a49, a65
+; CHECK-NEXT: v_accvgpr_mov_b32 a50, a66
+; CHECK-NEXT: v_accvgpr_mov_b32 a51, a67
+; CHECK-NEXT: v_accvgpr_mov_b32 a52, a68
+; CHECK-NEXT: v_accvgpr_mov_b32 a53, a69
+; CHECK-NEXT: v_accvgpr_mov_b32 a54, a70
+; CHECK-NEXT: v_accvgpr_mov_b32 a55, a71
+; CHECK-NEXT: v_accvgpr_mov_b32 a56, a72
+; CHECK-NEXT: v_accvgpr_mov_b32 a57, a73
+; CHECK-NEXT: v_accvgpr_mov_b32 a58, a74
+; CHECK-NEXT: v_accvgpr_mov_b32 a59, a75
+; CHECK-NEXT: v_accvgpr_mov_b32 a60, a76
+; CHECK-NEXT: v_accvgpr_mov_b32 a61, a77
+; CHECK-NEXT: v_accvgpr_mov_b32 a62, a78
+; CHECK-NEXT: v_accvgpr_mov_b32 a63, a79
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[4:19], v[30:31], v[30:31], v[4:19]
+; CHECK-NEXT: v_accvgpr_mov_b32 a64, a80
+; CHECK-NEXT: v_accvgpr_mov_b32 a65, a81
+; CHECK-NEXT: v_accvgpr_mov_b32 a66, a82
+; CHECK-NEXT: v_accvgpr_mov_b32 a67, a83
+; CHECK-NEXT: v_accvgpr_mov_b32 a68, a84
+; CHECK-NEXT: v_accvgpr_mov_b32 a69, a85
+; CHECK-NEXT: v_accvgpr_mov_b32 a70, a86
+; CHECK-NEXT: v_accvgpr_mov_b32 a71, a87
+; CHECK-NEXT: v_accvgpr_mov_b32 a72, a88
+; CHECK-NEXT: v_accvgpr_mov_b32 a73, a89
+; CHECK-NEXT: v_accvgpr_mov_b32 a74, a90
+; CHECK-NEXT: v_accvgpr_mov_b32 a75, a91
+; CHECK-NEXT: v_accvgpr_mov_b32 a76, a92
+; CHECK-NEXT: v_accvgpr_mov_b32 a77, a93
+; CHECK-NEXT: v_accvgpr_mov_b32 a78, a94
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[48:63], v[30:31], v[30:31], a[48:63]
+; CHECK-NEXT: v_accvgpr_mov_b32 a79, a95
+; CHECK-NEXT: v_pk_mov_b32 v[38:39], v[220:221], v[220:221] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[40:41], v[222:223], v[222:223] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[42:43], v[224:225], v[224:225] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[44:45], v[226:227], v[226:227] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[46:47], v[228:229], v[228:229] op_sel:[0,1]
+; CHECK-NEXT: v_mov_b32_e32 v130, v3
+; CHECK-NEXT: v_mov_b32_e32 v131, v3
+; CHECK-NEXT: v_mov_b32_e32 v132, v3
+; CHECK-NEXT: v_mov_b32_e32 v133, v3
+; CHECK-NEXT: v_mov_b32_e32 v134, v3
+; CHECK-NEXT: v_mov_b32_e32 v135, v3
+; CHECK-NEXT: v_mov_b32_e32 v136, v3
+; CHECK-NEXT: v_mov_b32_e32 v137, v3
+; CHECK-NEXT: v_mov_b32_e32 v138, v3
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[64:79], v[30:31], v[30:31], a[64:79]
+; CHECK-NEXT: v_mov_b32_e32 v139, v3
+; CHECK-NEXT: v_mov_b32_e32 v140, v3
+; CHECK-NEXT: v_mov_b32_e32 v141, v3
+; CHECK-NEXT: v_mov_b32_e32 v142, v3
+; CHECK-NEXT: v_mov_b32_e32 v143, v3
+; CHECK-NEXT: v_mov_b32_e32 v144, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a2, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a3, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a4, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a5, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a6, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a7, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a8, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a9, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a10, v3
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[86:101], v[30:31], v[30:31], v[86:101]
+; CHECK-NEXT: v_accvgpr_write_b32 a11, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a12, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a13, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a14, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a15, v3
+; CHECK-NEXT: v_mov_b32_e32 v163, v3
+; CHECK-NEXT: v_mov_b32_e32 v164, v3
+; CHECK-NEXT: v_mov_b32_e32 v165, v3
+; CHECK-NEXT: v_mov_b32_e32 v166, v3
+; CHECK-NEXT: v_mov_b32_e32 v167, v3
+; CHECK-NEXT: v_mov_b32_e32 v168, v3
+; CHECK-NEXT: v_mov_b32_e32 v169, v3
+; CHECK-NEXT: v_mov_b32_e32 v170, v3
+; CHECK-NEXT: v_mov_b32_e32 v171, v3
+; CHECK-NEXT: v_mov_b32_e32 v172, v3
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[32:47], v[30:31], v[30:31], v[32:47]
+; CHECK-NEXT: v_mov_b32_e32 v173, v3
+; CHECK-NEXT: v_mov_b32_e32 v174, v3
+; CHECK-NEXT: v_mov_b32_e32 v175, v3
+; CHECK-NEXT: v_mov_b32_e32 v176, v3
+; CHECK-NEXT: v_mov_b32_e32 v177, v3
+; CHECK-NEXT: v_mov_b32_e32 v147, v3
+; CHECK-NEXT: v_mov_b32_e32 v148, v3
+; CHECK-NEXT: v_mov_b32_e32 v149, v3
+; CHECK-NEXT: v_mov_b32_e32 v150, v3
+; CHECK-NEXT: v_mov_b32_e32 v151, v3
+; CHECK-NEXT: v_mov_b32_e32 v152, v3
+; CHECK-NEXT: v_mov_b32_e32 v153, v3
+; CHECK-NEXT: v_mov_b32_e32 v154, v3
+; CHECK-NEXT: v_mov_b32_e32 v155, v3
+; CHECK-NEXT: v_mov_b32_e32 v156, v3
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[4:19], v[30:31], v[30:31], v[4:19]
+; CHECK-NEXT: v_mov_b32_e32 v157, v3
+; CHECK-NEXT: v_mov_b32_e32 v158, v3
+; CHECK-NEXT: v_mov_b32_e32 v159, v3
+; CHECK-NEXT: v_mov_b32_e32 v160, v3
+; CHECK-NEXT: v_mov_b32_e32 v161, v3
+; CHECK-NEXT: v_accvgpr_read_b32 v68, a64
+; CHECK-NEXT: v_mov_b32_e32 v51, v3
+; CHECK-NEXT: v_mov_b32_e32 v52, v3
+; CHECK-NEXT: v_mov_b32_e32 v53, v3
+; CHECK-NEXT: v_mov_b32_e32 v54, v3
+; CHECK-NEXT: v_mov_b32_e32 v55, v3
+; CHECK-NEXT: v_mov_b32_e32 v56, v3
+; CHECK-NEXT: v_mov_b32_e32 v57, v3
+; CHECK-NEXT: v_mov_b32_e32 v58, v3
+; CHECK-NEXT: v_mov_b32_e32 v59, v3
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[32:47], v[30:31], v[30:31], a[32:47]
+; CHECK-NEXT: v_mov_b32_e32 v60, v3
+; CHECK-NEXT: v_mov_b32_e32 v61, v3
+; CHECK-NEXT: v_mov_b32_e32 v62, v3
+; CHECK-NEXT: v_mov_b32_e32 v63, v3
+; CHECK-NEXT: v_mov_b32_e32 v64, v3
+; CHECK-NEXT: v_mov_b32_e32 v65, v3
+; CHECK-NEXT: v_mov_b32_e32 v197, v3
+; CHECK-NEXT: v_mov_b32_e32 v198, v3
+; CHECK-NEXT: v_mov_b32_e32 v199, v3
+; CHECK-NEXT: v_mov_b32_e32 v200, v3
+; CHECK-NEXT: v_mov_b32_e32 v201, v3
+; CHECK-NEXT: v_mov_b32_e32 v202, v3
+; CHECK-NEXT: v_mov_b32_e32 v203, v3
+; CHECK-NEXT: v_mov_b32_e32 v204, v3
+; CHECK-NEXT: v_mov_b32_e32 v205, v3
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[48:63], v[30:31], v[30:31], a[48:63]
+; CHECK-NEXT: v_mov_b32_e32 v206, v3
+; CHECK-NEXT: v_mov_b32_e32 v207, v3
+; CHECK-NEXT: v_mov_b32_e32 v208, v3
+; CHECK-NEXT: v_mov_b32_e32 v209, v3
+; CHECK-NEXT: v_mov_b32_e32 v210, v3
+; CHECK-NEXT: v_mov_b32_e32 v211, v3
+; CHECK-NEXT: v_pk_mov_b32 v[228:229], v[46:47], v[46:47] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[226:227], v[44:45], v[44:45] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[224:225], v[42:43], v[42:43] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[222:223], v[40:41], v[40:41] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[220:221], v[38:39], v[38:39] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[218:219], v[36:37], v[36:37] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[216:217], v[34:35], v[34:35] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[214:215], v[32:33], v[32:33] op_sel:[0,1]
+; CHECK-NEXT: v_accvgpr_mov_b32 a18, a32
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[112:127], v[30:31], v[30:31], a[112:127]
+; CHECK-NEXT: v_accvgpr_mov_b32 a19, a33
+; CHECK-NEXT: v_accvgpr_mov_b32 a20, a34
+; CHECK-NEXT: v_accvgpr_mov_b32 a21, a35
+; CHECK-NEXT: v_accvgpr_read_b32 v32, a48
+; CHECK-NEXT: v_accvgpr_mov_b32 a22, a36
+; CHECK-NEXT: v_accvgpr_mov_b32 a23, a37
+; CHECK-NEXT: v_accvgpr_mov_b32 a24, a38
+; CHECK-NEXT: v_accvgpr_mov_b32 a25, a39
+; CHECK-NEXT: v_accvgpr_mov_b32 a26, a40
+; CHECK-NEXT: v_accvgpr_mov_b32 a27, a41
+; CHECK-NEXT: v_accvgpr_mov_b32 a28, a42
+; CHECK-NEXT: v_accvgpr_mov_b32 a29, a43
+; CHECK-NEXT: v_accvgpr_mov_b32 a30, a44
+; CHECK-NEXT: v_accvgpr_mov_b32 a31, a45
+; CHECK-NEXT: v_accvgpr_mov_b32 a32, a46
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[130:145], v[30:31], v[30:31], v[130:145]
+; CHECK-NEXT: v_accvgpr_mov_b32 a33, a47
+; CHECK-NEXT: v_pk_mov_b32 v[20:21], v[222:223], v[222:223] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[22:23], v[224:225], v[224:225] op_sel:[0,1]
+; CHECK-NEXT: v_accvgpr_mov_b32 a48, a112
+; CHECK-NEXT: v_accvgpr_mov_b32 a49, a113
+; CHECK-NEXT: v_accvgpr_mov_b32 a50, a114
+; CHECK-NEXT: v_accvgpr_mov_b32 a51, a115
+; CHECK-NEXT: v_accvgpr_mov_b32 a52, a116
+; CHECK-NEXT: v_accvgpr_mov_b32 a53, a117
+; CHECK-NEXT: v_accvgpr_mov_b32 a54, a118
+; CHECK-NEXT: v_accvgpr_mov_b32 a55, a119
+; CHECK-NEXT: v_accvgpr_mov_b32 a56, a120
+; CHECK-NEXT: v_accvgpr_mov_b32 a57, a121
+; CHECK-NEXT: v_accvgpr_mov_b32 a58, a122
+; CHECK-NEXT: v_accvgpr_mov_b32 a59, a123
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[96:111], v[30:31], v[30:31], a[96:111]
+; CHECK-NEXT: v_accvgpr_mov_b32 a60, a124
+; CHECK-NEXT: v_accvgpr_mov_b32 a61, a125
+; CHECK-NEXT: v_accvgpr_mov_b32 a62, a126
+; CHECK-NEXT: v_accvgpr_mov_b32 a63, a127
+; CHECK-NEXT: v_pk_mov_b32 v[24:25], v[226:227], v[226:227] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[26:27], v[228:229], v[228:229] op_sel:[0,1]
+; CHECK-NEXT: v_mov_b32_e32 v115, v3
+; CHECK-NEXT: v_mov_b32_e32 v116, v3
+; CHECK-NEXT: v_mov_b32_e32 v117, v3
+; CHECK-NEXT: v_mov_b32_e32 v118, v3
+; CHECK-NEXT: v_mov_b32_e32 v119, v3
+; CHECK-NEXT: v_mov_b32_e32 v120, v3
+; CHECK-NEXT: v_mov_b32_e32 v121, v3
+; CHECK-NEXT: v_mov_b32_e32 v122, v3
+; CHECK-NEXT: v_mov_b32_e32 v123, v3
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[80:95], v[30:31], v[30:31], a[64:79]
+; CHECK-NEXT: v_mov_b32_e32 v124, v3
+; CHECK-NEXT: v_mov_b32_e32 v125, v3
+; CHECK-NEXT: v_mov_b32_e32 v126, v3
+; CHECK-NEXT: v_mov_b32_e32 v127, v3
+; CHECK-NEXT: v_mov_b32_e32 v128, v3
+; CHECK-NEXT: v_mov_b32_e32 v129, v3
+; CHECK-NEXT: v_pk_mov_b32 v[34:35], v[214:215], v[214:215] op_sel:[0,1]
+; CHECK-NEXT: v_mov_b32_e32 v67, v66
+; CHECK-NEXT: v_mov_b32_e32 v243, v242
+; CHECK-NEXT: v_accvgpr_mov_b32 a17, a16
+; CHECK-NEXT: v_pk_mov_b32 v[36:37], v[216:217], v[216:217] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[38:39], v[218:219], v[218:219] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[40:41], v[220:221], v[220:221] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[42:43], v[222:223], v[222:223] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[44:45], v[224:225], v[224:225] op_sel:[0,1]
+; CHECK-NEXT: v_accvgpr_mov_b32 a79, a15
+; CHECK-NEXT: v_accvgpr_mov_b32 a78, a14
+; CHECK-NEXT: v_accvgpr_mov_b32 a77, a13
+; CHECK-NEXT: v_accvgpr_mov_b32 a76, a12
+; CHECK-NEXT: v_accvgpr_mov_b32 a75, a11
+; CHECK-NEXT: v_accvgpr_mov_b32 a74, a10
+; CHECK-NEXT: v_accvgpr_mov_b32 a73, a9
+; CHECK-NEXT: v_accvgpr_mov_b32 a72, a8
+; CHECK-NEXT: v_accvgpr_mov_b32 a71, a7
+; CHECK-NEXT: v_accvgpr_mov_b32 a70, a6
+; CHECK-NEXT: v_accvgpr_mov_b32 a69, a5
+; CHECK-NEXT: v_accvgpr_mov_b32 a68, a4
+; CHECK-NEXT: v_accvgpr_mov_b32 a67, a3
+; CHECK-NEXT: v_accvgpr_mov_b32 a66, a2
+; CHECK-NEXT: v_accvgpr_mov_b32 a65, a1
+; CHECK-NEXT: v_accvgpr_mov_b32 a64, a0
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[146:161], v[30:31], v[30:31], v[146:161]
+; CHECK-NEXT: v_accvgpr_mov_b32 a1, a0
+; CHECK-NEXT: v_pk_mov_b32 v[46:47], v[226:227], v[226:227] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[48:49], v[228:229], v[228:229] op_sel:[0,1]
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[64:79], v[30:31], v[30:31], a[64:79]
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[86:101], v[30:31], v[30:31], v[86:101]
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[162:177], v[30:31], v[30:31], v[162:177]
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[50:65], v[30:31], v[30:31], v[50:65]
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[196:211], v[30:31], v[30:31], v[196:211]
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[4:19], v[30:31], v[28:29], v[4:19]
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[130:145], v[30:31], v[30:31], v[130:145]
+; CHECK-NEXT: s_nop 15
+; CHECK-NEXT: s_nop 1
+; CHECK-NEXT: v_pk_mov_b32 v[12:13], v[214:215], v[214:215] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[14:15], v[216:217], v[216:217] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[16:17], v[218:219], v[218:219] op_sel:[0,1]
+; CHECK-NEXT: v_pk_mov_b32 v[18:19], v[220:221], v[220:221] op_sel:[0,1]
+; CHECK-NEXT: v_mov_b32_e32 v66, v11
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[96:111], v[30:31], v[30:31], a[96:111]
+; CHECK-NEXT: v_mov_b32_e32 v145, v130
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[80:95], v[30:31], v[30:31], a[80:95]
+; CHECK-NEXT: s_nop 15
+; CHECK-NEXT: s_nop 0
+; CHECK-NEXT: v_accvgpr_read_b32 v2, a96
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[18:33], v[30:31], v[30:31], a[18:33]
+; CHECK-NEXT: v_accvgpr_read_b32 v70, a80
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[48:63], v[30:31], v[30:31], a[48:63]
+; CHECK-NEXT: s_nop 15
+; CHECK-NEXT: s_nop 0
+; CHECK-NEXT: v_accvgpr_read_b32 v212, a18
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[12:27], v[30:31], v[30:31], v[12:27]
+; CHECK-NEXT: v_accvgpr_mov_b32 a16, a48
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[146:161], v[30:31], v[30:31], v[146:161]
+; CHECK-NEXT: s_nop 15
+; CHECK-NEXT: s_nop 0
+; CHECK-NEXT: v_mov_b32_e32 v35, v12
+; CHECK-NEXT: v_mov_b32_e32 v20, v68
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[86:101], v[30:31], v[30:31], v[86:101]
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 a[64:79], v[30:31], v[30:31], a[64:79]
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[162:177], v[30:31], v[30:31], v[162:177]
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[50:65], v[30:31], v[30:31], v[50:65]
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[196:211], v[30:31], v[30:31], v[196:211]
+; CHECK-NEXT: v_mfma_f32_32x32x8f16 v[114:129], v[30:31], v[30:31], v[114:129]
+; CHECK-NEXT: s_mov_b64 vcc, vcc
+; CHECK-NEXT: s_cbranch_vccnz .LBB0_1
+; CHECK-NEXT: ; %bb.2: ; %DummyReturnBlock
+; CHECK-NEXT: s_endpgm
+entry:
+ br label %do.body
+
+do.body:
+ %c_block_tile.sroa.37.0 = phi float [ 0.000000e+00, %entry ], [ %c_block_tile.sroa.70.0, %do.body ]
+ %c_block_tile.sroa.70.0 = phi float [ 0.000000e+00, %entry ], [ %ext0, %do.body ]
+ %c_block_tile.sroa.961.0 = phi float [ 0.000000e+00, %entry ], [ %ext1, %do.body ]
+ %c_block_tile.sroa.994.0 = phi float [ 0.000000e+00, %entry ], [ %ext2, %do.body ]
+ %c_block_tile.sroa.1060.0 = phi float [ 0.000000e+00, %entry ], [ %ext3, %do.body ]
+ %c_block_tile.sroa.1093.0 = phi float [ 0.000000e+00, %entry ], [ %c_block_tile.sroa.1060.0, %do.body ]
+ %c_block_tile.sroa.1588.0 = phi float [ 0.000000e+00, %entry ], [ %ext4, %do.body ]
+ %c_block_tile.sroa.1687.0 = phi float [ 0.000000e+00, %entry ], [ %ext5, %do.body ]
+ %c_block_tile.sroa.2578.0 = phi float [ 0.000000e+00, %entry ], [ %ext6, %do.body ]
+ %c_block_tile.sroa.2611.0 = phi float [ 0.000000e+00, %entry ], [ %ext7, %do.body ]
+ %c_block_tile.sroa.2644.0 = phi float [ 0.000000e+00, %entry ], [ %ext8, %do.body ]
+ %c_block_tile.sroa.2677.0 = phi float [ 0.000000e+00, %entry ], [ %ext9, %do.body ]
+ %c_block_tile.sroa.3568.0 = phi float [ 0.000000e+00, %entry ], [ %ext10, %do.body ]
+ %c_block_tile.sroa.3634.0 = phi float [ 0.000000e+00, %entry ], [ %ext11, %do.body ]
+ %c_block_tile.sroa.3898.0 = phi float [ 0.000000e+00, %entry ], [ %ext12, %do.body ]
+ %c_block_tile.sroa.3931.0 = phi float [ 0.000000e+00, %entry ], [ %ext13, %do.body ]
+ %c_block_tile.sroa.4624.0 = phi float [ 0.000000e+00, %entry ], [ %ext14, %do.body ]
+ %c_block_tile.sroa.4657.0 = phi float [ 0.000000e+00, %entry ], [ %ext15, %do.body ]
+ %c_block_tile.sroa.5185.0 = phi float [ 0.000000e+00, %entry ], [ %ext16, %do.body ]
+ %c_block_tile.sroa.5218.0 = phi float [ 0.000000e+00, %entry ], [ %ext17, %do.body ]
+ %c_block_tile.sroa.5746.0 = phi float [ 0.000000e+00, %entry ], [ %ext18, %do.body ]
+ %c_block_tile.sroa.5779.0 = phi float [ 0.000000e+00, %entry ], [ %c_block_tile.sroa.5746.0, %do.body ]
+ %c_block_tile.sroa.5812.0 = phi float [ 0.000000e+00, %entry ], [ %ext19, %do.body ]
+ %c_block_tile.sroa.6373.0 = phi float [ 0.000000e+00, %entry ], [ %c_block_tile.sroa.6406.0, %do.body ]
+ %c_block_tile.sroa.6406.0 = phi float [ 0.000000e+00, %entry ], [ %ext20, %do.body ]
+ %c_block_tile.sroa.7297.0 = phi float [ 0.000000e+00, %entry ], [ %ext21, %do.body ]
+ %c_block_tile.sroa.7363.0 = phi float [ 0.000000e+00, %entry ], [ %ext22, %do.body ]
+ %c_block_tile.sroa.7396.0 = phi float [ 0.000000e+00, %entry ], [ %ext23, %do.body ]
+ %c_block_tile.sroa.7429.0 = phi float [ 0.000000e+00, %entry ], [ %ext24, %do.body ]
+ %v0 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.37.0, i64 0
+ %v1 = insertelement <16 x float> %v0, float %c_block_tile.sroa.70.0, i64 0
+ %0 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v1, i32 0, i32 0, i32 0)
+ %v2 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.961.0, i64 13
+ %v3 = insertelement <16 x float> %v2, float %c_block_tile.sroa.994.0, i64 14
+ %v4 = insertelement <16 x float> %v3, float +qnan, i64 0
+ %1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v4, i32 0, i32 0, i32 0)
+ %v5 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.1588.0, i64 0
+ %v6 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.1687.0, i64 3
+ %v7 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.2578.0, i64 0
+ %v8 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.2611.0, i64 0
+ %v9 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.2677.0, i64 0
+ %v10 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.3568.0, i64 0
+ %v11 = insertelement <16 x float> splat (float 1.000000e+00), float %c_block_tile.sroa.3634.0, i64 0
+ %2 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v11, i32 0, i32 0, i32 0)
+ %v12 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.3898.0, i64 1
+ %v13 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.4624.0, i64 0
+ %v14 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.5185.0, i64 0
+ %v15 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.5218.0, i64 0
+ %v16 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.5746.0, i64 14
+ %v17 = insertelement <16 x float> %v16, float %c_block_tile.sroa.5779.0, i64 15
+ %3 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v17, i32 0, i32 0, i32 0)
+ %v18 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.5812.0, i64 0
+ %v19 = insertelement <16 x float> %v18, float %c_block_tile.sroa.5812.0, i64 0
+ %v20 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.6373.0, i64 1
+ %v21 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.7297.0, i64 0
+ %v22 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.7363.0, i64 15
+ %v23 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.7396.0, i64 0
+ %v24 = insertelement <16 x float> %v23, float %c_block_tile.sroa.7429.0, i64 1
+ %4 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v15, i32 0, i32 0, i32 0)
+ %5 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %4, i32 0, i32 0, i32 0)
+ %6 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v22, i32 0, i32 0, i32 0)
+ %7 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %6, i32 0, i32 0, i32 0)
+ %8 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v6, i32 0, i32 0, i32 0)
+ %9 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %0, i32 0, i32 0, i32 0)
+ %10 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %2, i32 0, i32 0, i32 0)
+ %v25 = insertelement <16 x float> %v12, float %c_block_tile.sroa.3931.0, i64 0
+ %11 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v25, i32 0, i32 0, i32 0)
+ %v26 = insertelement <16 x float> %v20, float %c_block_tile.sroa.6406.0, i64 0
+ %12 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v26, i32 0, i32 0, i32 0)
+ %13 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v24, i32 0, i32 0, i32 0)
+ %14 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %9, i32 0, i32 0, i32 0)
+ %15 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %11, i32 0, i32 0, i32 0)
+ %16 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %3, i32 0, i32 0, i32 0)
+ %17 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %1, i32 0, i32 0, i32 0)
+ %18 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %17, i32 0, i32 0, i32 0)
+ %v27 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.1060.0, i64 0
+ %v28 = insertelement <16 x float> %v27, float %c_block_tile.sroa.1093.0, i64 1
+ %19 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v28, i32 0, i32 0, i32 0)
+ %20 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %8, i32 0, i32 0, i32 0)
+ %21 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %20, i32 0, i32 0, i32 0)
+ %22 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %16, i32 0, i32 0, i32 0)
+ %23 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %12, i32 0, i32 0, i32 0)
+ %ext0 = extractelement <16 x float> %14, i64 0
+ %24 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %18, i32 0, i32 0, i32 0)
+ %ext1 = extractelement <16 x float> %24, i64 0
+ %ext2 = extractelement <16 x float> %17, i64 0
+ %25 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %19, i32 0, i32 0, i32 0)
+ %ext3 = extractelement <16 x float> %25, i64 0
+ %ext4 = extractelement <16 x float> %8, i64 0
+ %ext5 = extractelement <16 x float> %21, i64 0
+ %26 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v8, i32 0, i32 0, i32 0)
+ %ext6 = extractelement <16 x float> %26, i64 0
+ %27 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %26, i32 0, i32 0, i32 0)
+ %ext7 = extractelement <16 x float> %27, i64 0
+ %v29 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.2644.0, i64 0
+ %28 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v29, i32 0, i32 0, i32 0)
+ %29 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %28, i32 0, i32 0, i32 0)
+ %ext8 = extractelement <16 x float> %29, i64 0
+ %ext9 = extractelement <16 x float> %28, i64 0
+ %ext10 = extractelement <16 x float> %2, i64 0
+ %30 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %10, i32 0, i32 0, i32 0)
+ %ext11 = extractelement <16 x float> %30, i64 0
+ %ext12 = extractelement <16 x float> %11, i64 0
+ %31 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %15, i32 0, i32 0, i32 0)
+ %32 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %31, i32 0, i32 0, i32 0)
+ %33 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> splat (half 0xH3C00), <16 x float> %32, i32 0, i32 0, i32 0)
+ %ext13 = extractelement <16 x float> %33, i64 7
+ %v30 = insertelement <16 x float> zeroinitializer, float %c_block_tile.sroa.4657.0, i64 0
+ %34 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v30, i32 0, i32 0, i32 0)
+ %ext14 = extractelement <16 x float> %34, i64 0
+ %35 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %34, i32 0, i32 0, i32 0)
+ %ext15 = extractelement <16 x float> %35, i64 0
+ %ext16 = extractelement <16 x float> %4, i64 0
+ %ext17 = extractelement <16 x float> %5, i64 0
+ %ext18 = extractelement <16 x float> %22, i64 0
+ %36 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %v19, i32 0, i32 0, i32 0)
+ %ext19 = extractelement <16 x float> %36, i64 0
+ %37 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %23, i32 0, i32 0, i32 0)
+ %ext20 = extractelement <16 x float> %37, i64 0
+ %ext21 = extractelement <16 x float> %6, i64 0
+ %ext22 = extractelement <16 x float> %7, i64 0
+ %ext23 = extractelement <16 x float> %13, i64 0
+ %38 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <16 x float> %13, i32 0, i32 0, i32 0)
+ %ext24 = extractelement <16 x float> %38, i64 0
+ br label %do.body
+}
+
+declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half>, <4 x half>, <16 x float>, i32 immarg, i32 immarg, i32 immarg)
+
+attributes #0 = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-lds-size"="32768" }
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
index 94e9891652694..21dd9ff425623 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store.ll
@@ -1,6 +1,6 @@
; REQUIRES: asserts
; RUN: llc -mtriple=amdgpu9.0a-amd-amdhsa -O3 \
-; RUN: -amdgpu-use-amdgpu-trackers=1 -verify-machineinstrs \
+; RUN: -amdgpu-use-amdgpu-trackers=1 -vgpr-regalloc=basic -verify-machineinstrs \
; RUN: -stop-after=amdgpu-rewrite-agpr-copy-mfma \
; RUN: -debug-only=amdgpu-rewrite-agpr-copy-mfma -filetype=null %s 2>&1 | FileCheck %s
@@ -145,4 +145,4 @@ do.body:
declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half>, <4 x half>, <16 x float>, i32 immarg, i32 immarg, i32 immarg)
-attributes #0 = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-lds-size"="32768" }
+attributes #0 = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-lds-size"="32768" "amdgpu-waves-per-eu"="2" }
More information about the llvm-commits
mailing list