[AMDGPU] Autogen checks for mfma-no-register-aliasing.ll

[jrbyrnes]

For an upcoming RegisterCoalescer PR

[llvmbot]

@llvm/pr-subscribers-backend-amdgpu

Author: Jeffrey Byrnes (jrbyrnes)

Changes

For an upcoming RegisterCoalescer PR

---

Patch is 52.70 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/132117.diff

1 Files Affected:

• (modified) llvm/test/CodeGen/AMDGPU/mfma-no-register-aliasing.ll (+1009-31)

diff --git a/llvm/test/CodeGen/AMDGPU/mfma-no-register-aliasing.ll b/llvm/test/CodeGen/AMDGPU/mfma-no-register-aliasing.ll
index f76580b94e13c..21af2dde2c4bf 100644
--- a/llvm/test/CodeGen/AMDGPU/mfma-no-register-aliasing.ll
+++ b/llvm/test/CodeGen/AMDGPU/mfma-no-register-aliasing.ll
@@ -1,9 +1,10 @@
-; RUN: llc -mtriple=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GREEDY,GREEDY908 %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GREEDY,GREEDY90A %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -early-live-intervals -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GREEDY,GREEDY90A %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GREEDY,GREEDY90A %s
-; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GREEDY,GREEDY90A-GISEL %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -sgpr-regalloc=fast -vgpr-regalloc=fast -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN,FAST %s
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc -mtriple=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GREEDY908 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GREEDY90A %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -early-live-intervals -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GREEDY90A %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GREEDY942 %s
+; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GREEDY90A-GISEL %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -sgpr-regalloc=fast -vgpr-regalloc=fast -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=FAST90A %s
 
 ; This is better with 90a
 
@@ -14,13 +15,612 @@ declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>
 declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float, float, <16 x float>, i32, i32, i32)
 declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32, i32, i32)
 
-; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32:
-; GREEDY: v_mfma_f32_32x32x1{{.*}} a[0:31], v{{[0-9]+}}, v{{[0-9]+}}, a[0:31]
-; GREEDY: v_mfma_f32_32x32x1{{.*}} a[32:63], v{{[0-9]+}}, v{{[0-9]+}}, a[0:31]
-; FAST:   v_mfma_f32_32x32x1{{.*}} a[32:63], v{{[0-9]+}}, v{{[0-9]+}}, a[32:63]
-; FAST:   v_mfma_f32_32x32x1{{.*}} a[0:31], v{{[0-9]+}}, v{{[0-9]+}}, a[32:63]
-; GCN:    v_mfma_f32_32x32x1{{.*}} a[0:31], v{{[0-9]+}}, v{{[0-9]+}}, a[0:31]
 define amdgpu_kernel void @test_mfma_f32_32x32x1f32(ptr addrspace(1) %arg) #0 {
+; GREEDY908-LABEL: test_mfma_f32_32x32x1f32:
+; GREEDY908:       ; %bb.0: ; %bb
+; GREEDY908-NEXT:    s_load_dwordx2 s[34:35], s[4:5], 0x24
+; GREEDY908-NEXT:    s_waitcnt lgkmcnt(0)
+; GREEDY908-NEXT:    s_load_dwordx16 s[16:31], s[34:35], 0x0
+; GREEDY908-NEXT:    s_load_dwordx16 s[0:15], s[34:35], 0x40
+; GREEDY908-NEXT:    s_waitcnt lgkmcnt(0)
+; GREEDY908-NEXT:    v_mov_b32_e32 v0, s16
+; GREEDY908-NEXT:    v_mov_b32_e32 v1, s17
+; GREEDY908-NEXT:    v_mov_b32_e32 v2, s18
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a0, v0
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a1, v1
+; GREEDY908-NEXT:    v_mov_b32_e32 v0, s22
+; GREEDY908-NEXT:    v_mov_b32_e32 v1, s23
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a2, v2
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a6, v0
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a7, v1
+; GREEDY908-NEXT:    v_mov_b32_e32 v0, s24
+; GREEDY908-NEXT:    v_mov_b32_e32 v1, s25
+; GREEDY908-NEXT:    v_mov_b32_e32 v2, s26
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a8, v0
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a9, v1
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a10, v2
+; GREEDY908-NEXT:    v_mov_b32_e32 v0, s27
+; GREEDY908-NEXT:    v_mov_b32_e32 v1, s28
+; GREEDY908-NEXT:    v_mov_b32_e32 v2, s29
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a11, v0
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a12, v1
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a13, v2
+; GREEDY908-NEXT:    v_mov_b32_e32 v0, s30
+; GREEDY908-NEXT:    v_mov_b32_e32 v1, s31
+; GREEDY908-NEXT:    v_mov_b32_e32 v2, s0
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a14, v0
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a15, v1
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a16, v2
+; GREEDY908-NEXT:    v_mov_b32_e32 v0, s1
+; GREEDY908-NEXT:    v_mov_b32_e32 v1, s2
+; GREEDY908-NEXT:    v_mov_b32_e32 v2, s3
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a17, v0
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a18, v1
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a19, v2
+; GREEDY908-NEXT:    v_mov_b32_e32 v0, s4
+; GREEDY908-NEXT:    v_mov_b32_e32 v1, s5
+; GREEDY908-NEXT:    v_mov_b32_e32 v2, s6
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a20, v0
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a21, v1
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a22, v2
+; GREEDY908-NEXT:    v_mov_b32_e32 v0, s7
+; GREEDY908-NEXT:    v_mov_b32_e32 v1, s8
+; GREEDY908-NEXT:    v_mov_b32_e32 v2, s9
+; GREEDY908-NEXT:    v_mov_b32_e32 v3, s19
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a23, v0
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a24, v1
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a25, v2
+; GREEDY908-NEXT:    v_mov_b32_e32 v0, s10
+; GREEDY908-NEXT:    v_mov_b32_e32 v1, s11
+; GREEDY908-NEXT:    v_mov_b32_e32 v2, s12
+; GREEDY908-NEXT:    v_mov_b32_e32 v4, s20
+; GREEDY908-NEXT:    v_mov_b32_e32 v5, s21
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a3, v3
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a26, v0
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a27, v1
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a28, v2
+; GREEDY908-NEXT:    v_mov_b32_e32 v0, s13
+; GREEDY908-NEXT:    v_mov_b32_e32 v1, s14
+; GREEDY908-NEXT:    v_mov_b32_e32 v2, s15
+; GREEDY908-NEXT:    v_mov_b32_e32 v3, 1.0
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a4, v4
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a5, v5
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a29, v0
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a30, v1
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a31, v2
+; GREEDY908-NEXT:    v_mov_b32_e32 v0, 2.0
+; GREEDY908-NEXT:    v_mov_b32_e32 v4, 0
+; GREEDY908-NEXT:    s_nop 0
+; GREEDY908-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v3, v0, a[0:31]
+; GREEDY908-NEXT:    v_mfma_f32_32x32x1f32 a[32:63], v3, v0, a[0:31]
+; GREEDY908-NEXT:    s_nop 7
+; GREEDY908-NEXT:    s_nop 7
+; GREEDY908-NEXT:    s_nop 1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a32
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v5, a61
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v6, a60
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a2, v1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a33
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v7, a59
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v8, a58
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a3, v1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a34
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v9, a57
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v10, a56
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a4, v1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a35
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v11, a55
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v12, a54
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a5, v1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a36
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v13, a53
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v14, a52
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a6, v1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a37
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v15, a51
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v16, a50
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a7, v1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a38
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v17, a49
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v18, a48
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a8, v1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a39
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v19, a47
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v2, a46
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a9, v1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a40
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a16, v2
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a17, v19
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a10, v1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a41
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a18, v18
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a19, v17
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a11, v1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a42
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a20, v16
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a21, v15
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a12, v1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a43
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a22, v14
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a23, v13
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a13, v1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a44
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a24, v12
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a25, v11
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a14, v1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a45
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a26, v10
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a27, v9
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a15, v1
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a28, v8
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a29, v7
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a30, v6
+; GREEDY908-NEXT:    v_accvgpr_write_b32 a31, v5
+; GREEDY908-NEXT:    s_nop 0
+; GREEDY908-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v3, v0, a[0:31]
+; GREEDY908-NEXT:    s_nop 7
+; GREEDY908-NEXT:    s_nop 7
+; GREEDY908-NEXT:    s_nop 1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v3, a27
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v2, a26
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a25
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v0, a24
+; GREEDY908-NEXT:    s_nop 1
+; GREEDY908-NEXT:    global_store_dwordx4 v4, v[0:3], s[34:35] offset:96
+; GREEDY908-NEXT:    s_nop 0
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v3, a31
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v2, a30
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a29
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v0, a28
+; GREEDY908-NEXT:    s_nop 1
+; GREEDY908-NEXT:    global_store_dwordx4 v4, v[0:3], s[34:35] offset:112
+; GREEDY908-NEXT:    s_nop 0
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v3, a19
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v2, a18
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a17
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v0, a16
+; GREEDY908-NEXT:    s_nop 1
+; GREEDY908-NEXT:    global_store_dwordx4 v4, v[0:3], s[34:35] offset:64
+; GREEDY908-NEXT:    s_nop 0
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v3, a23
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v2, a22
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a21
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v0, a20
+; GREEDY908-NEXT:    s_nop 1
+; GREEDY908-NEXT:    global_store_dwordx4 v4, v[0:3], s[34:35] offset:80
+; GREEDY908-NEXT:    s_nop 0
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v3, a11
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v2, a10
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a9
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v0, a8
+; GREEDY908-NEXT:    s_nop 1
+; GREEDY908-NEXT:    global_store_dwordx4 v4, v[0:3], s[34:35] offset:32
+; GREEDY908-NEXT:    s_nop 0
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v3, a15
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v2, a14
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a13
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v0, a12
+; GREEDY908-NEXT:    s_nop 1
+; GREEDY908-NEXT:    global_store_dwordx4 v4, v[0:3], s[34:35] offset:48
+; GREEDY908-NEXT:    s_nop 0
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v3, a3
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v2, a2
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a1
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v0, a0
+; GREEDY908-NEXT:    s_nop 1
+; GREEDY908-NEXT:    global_store_dwordx4 v4, v[0:3], s[34:35]
+; GREEDY908-NEXT:    s_nop 0
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v3, a7
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v2, a6
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v1, a5
+; GREEDY908-NEXT:    v_accvgpr_read_b32 v0, a4
+; GREEDY908-NEXT:    s_nop 1
+; GREEDY908-NEXT:    global_store_dwordx4 v4, v[0:3], s[34:35] offset:16
+; GREEDY908-NEXT:    s_endpgm
+;
+; GREEDY90A-LABEL: test_mfma_f32_32x32x1f32:
+; GREEDY90A:       ; %bb.0: ; %bb
+; GREEDY90A-NEXT:    s_load_dwordx2 s[34:35], s[4:5], 0x24
+; GREEDY90A-NEXT:    v_mov_b32_e32 v0, 1.0
+; GREEDY90A-NEXT:    v_mov_b32_e32 v1, 2.0
+; GREEDY90A-NEXT:    v_mov_b32_e32 v2, 0
+; GREEDY90A-NEXT:    s_waitcnt lgkmcnt(0)
+; GREEDY90A-NEXT:    s_load_dwordx16 s[16:31], s[34:35], 0x0
+; GREEDY90A-NEXT:    s_load_dwordx16 s[0:15], s[34:35], 0x40
+; GREEDY90A-NEXT:    s_waitcnt lgkmcnt(0)
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a0, s16
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a1, s17
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a2, s18
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a3, s19
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a4, s20
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a5, s21
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a6, s22
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a7, s23
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a8, s24
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a9, s25
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a10, s26
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a11, s27
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a12, s28
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a13, s29
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a14, s30
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a15, s31
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a16, s0
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a17, s1
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a18, s2
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a19, s3
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a20, s4
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a21, s5
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a22, s6
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a23, s7
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a24, s8
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a25, s9
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a26, s10
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a27, s11
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a28, s12
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a29, s13
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a30, s14
+; GREEDY90A-NEXT:    v_accvgpr_write_b32 a31, s15
+; GREEDY90A-NEXT:    s_nop 1
+; GREEDY90A-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]
+; GREEDY90A-NEXT:    v_mfma_f32_32x32x1f32 a[32:63], v0, v1, a[0:31]
+; GREEDY90A-NEXT:    s_nop 7
+; GREEDY90A-NEXT:    s_nop 7
+; GREEDY90A-NEXT:    s_nop 2
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a2, a32
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a3, a33
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a4, a34
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a5, a35
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a6, a36
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a7, a37
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a8, a38
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a9, a39
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a10, a40
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a11, a41
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a12, a42
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a13, a43
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a14, a44
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a15, a45
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a16, a46
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a17, a47
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a18, a48
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a19, a49
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a20, a50
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a21, a51
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a22, a52
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a23, a53
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a24, a54
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a25, a55
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a26, a56
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a27, a57
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a28, a58
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a29, a59
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a30, a60
+; GREEDY90A-NEXT:    v_accvgpr_mov_b32 a31, a61
+; GREEDY90A-NEXT:    s_nop 1
+; GREEDY90A-NEXT:    v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[0:31]
+; GREEDY90A-NEXT:    s_nop 7
+; GREEDY90A-NEXT:    s_nop 7
+; GREEDY90A-NEXT:    s_nop 2
+; GREEDY90A-NEXT:    global_store_dwordx4 v2, a[24:27], s[34:35] offset:96
+; GREEDY90A-NEXT:    global_store_dwordx4 v2, a[28:31], s[34:35] offset:112
+; GREEDY90A-NEXT:    global_store_dwordx4 v2, a[16:19], s[34:35] offset:64
+; GREEDY90A-NEXT:    global_store_dwordx4 v2, a[20:23], s[34:35] offset:80
+; GREEDY90A-NEXT:    global_store_dwordx4 v2, a[8:11], s[34:35] offset:32
+; GREEDY90A-NEXT:    global_store_dwordx4 v2, a[12:15], s[34:35] offset:48
+; GREEDY90A-NEXT:    global_store_dwordx4 v2, a[0:3], s[34:35]
+; GREEDY90A-NEXT:    global_store_dwordx4 v2, a[4:7], s[34:35] offset:16
+; GREEDY90A-NEXT:    s_endpgm
+;
+; GREEDY942-LABEL: test_mfma_f32_32x32x1f32:
+; GREEDY942:       ; %bb.0: ; %bb
+; GREEDY942-NEXT:    s_load_dwordx2 s[34:35], s[4:5], 0x24
+; GREEDY942-NEXT:    v_mov_b32_e32 v0, 1.0
+; GREEDY942-NEXT:    v_mov_b32_e32 v1, 2.0
+; GREEDY942-NEXT:    v_mov_b32_e32 v2, 0
+; GREEDY942-NEXT:    s_waitcnt lgkmcnt(0)
+; GREEDY942-NEXT:    s_load_dwordx16 s[16:31], s[34:35], 0x0
+; GREEDY942-NEXT:    s_load_dwordx16 s[0:15], s[34:35], 0x40
+; GREEDY942-NEXT:    s_waitcnt lgkmcnt(0)
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a0, s16
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a1, s17
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a2, s18
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a3, s19
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a4, s20
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a5, s21
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a6, s22
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a7, s23
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a8, s24
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a9, s25
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a10, s26
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a11, s27
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a12, s28
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a13, s29
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a14, s30
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a15, s31
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a16, s0
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a17, s1
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a18, s2
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a19, s3
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a20, s4
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a21, s5
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a22, s6
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a23, s7
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a24, s8
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a25, s9
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a26, s10
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a27, s11
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a28, s12
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a29, s13
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a30, s14
+; GREEDY942-NEXT:    v_accvgpr_write_b32 a31, s15
+; GREEDY942-NEXT:    s_nop 1
+; GREEDY942-NEXT:    v_mfma_f32_32x32x1_2b_f32 a[0:31], v0, v1, a[0:31]
+; GREEDY942-NEXT:    v_mfma_f32_32x32x1_2b_f32 a[32:63], v0, v1, a[0:31]
+; GREEDY942-NEXT:    s_nop 7
+; GREEDY942-NEXT:    s_nop 7
+; GREEDY942-NEXT:    s_nop 1
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a2, a32
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a3, a33
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a4, a34
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a5, a35
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a6, a36
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a7, a37
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a8, a38
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a9, a39
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a10, a40
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a11, a41
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a12, a42
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a13, a43
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a14, a44
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a15, a45
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a16, a46
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a17, a47
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a18, a48
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a19, a49
+; GREEDY942-NEXT:    v_accvgpr_mov_b32 a20, a50
+; GREEDY942-NEXT:    v_accvgpr_mov_b...
[truncated]