AMDGPU: Convert some mfma tests to generated checks

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.i8.ll

@@ -1,41 +1,172 @@
-; RUN: llc -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX908 %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx908 -mattr=-mfma-inline-literal-bug < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX908 %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX90A %s
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck --check-prefixes=GCN,GFX908 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx908 -mattr=-mfma-inline-literal-bug < %s | FileCheck --check-prefixes=GCN,GFX908 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck --check-prefixes=GCN,GFX90A %s
 
 declare <16 x i32> @llvm.amdgcn.mfma.i32.32x32x8i8(i32, i32, <16 x i32>, i32, i32, i32)
 declare <4 x i32> @llvm.amdgcn.mfma.i32.16x16x16i8(i32, i32, <4 x i32>, i32, i32, i32)
 
-; GCN-LABEL: {{^}}test_mfma_i32_32x32x8i8:
-; GCN-DAG:         v_mov_b32_e32 [[TWO:v[0-9]+]], 2
-; GCN-DAG:         v_mov_b32_e32 [[ONE:v[0-9]+]], 1
-; GCN-DAG:         s_load_dwordx16
-; GFX908-DAG-COUNT-16: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
-; GFX90A-COUNT-16: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
-; GCN:             v_mfma_i32_32x32x8i8 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
-; GFX908-COUNT-16: v_accvgpr_read_b32
-; GFX908:          global_store_dwordx4
-; GFX90A-NOT:      v_accvgpr_read_b32
-; GFX90A-COUNT-4:  global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}]
 define amdgpu_kernel void @test_mfma_i32_32x32x8i8(ptr addrspace(1) %arg) #0 {
+; GFX908-LABEL: test_mfma_i32_32x32x8i8:
+; GFX908:       ; %bb.0: ; %bb
+; GFX908-NEXT:    s_load_dwordx2 s[16:17], s[4:5], 0x24
+; GFX908-NEXT:    v_mov_b32_e32 v0, 1
+; GFX908-NEXT:    v_mov_b32_e32 v16, 0
+; GFX908-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX908-NEXT:    s_load_dwordx16 s[0:15], s[16:17], 0x0
+; GFX908-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX908-NEXT:    v_mov_b32_e32 v17, s0
+; GFX908-NEXT:    v_mov_b32_e32 v1, s1
+; GFX908-NEXT:    v_mov_b32_e32 v2, s2
+; GFX908-NEXT:    v_accvgpr_write_b32 a0, v17
+; GFX908-NEXT:    v_mov_b32_e32 v17, s3
+; GFX908-NEXT:    v_accvgpr_write_b32 a1, v1
+; GFX908-NEXT:    v_accvgpr_write_b32 a2, v2
+; GFX908-NEXT:    v_accvgpr_write_b32 a3, v17
+; GFX908-NEXT:    v_mov_b32_e32 v1, s4
+; GFX908-NEXT:    v_mov_b32_e32 v2, s5
+; GFX908-NEXT:    v_mov_b32_e32 v17, s6
+; GFX908-NEXT:    v_accvgpr_write_b32 a4, v1
+; GFX908-NEXT:    v_accvgpr_write_b32 a5, v2
+; GFX908-NEXT:    v_accvgpr_write_b32 a6, v17
+; GFX908-NEXT:    v_mov_b32_e32 v1, s7
+; GFX908-NEXT:    v_mov_b32_e32 v2, s8
+; GFX908-NEXT:    v_mov_b32_e32 v17, s9
+; GFX908-NEXT:    v_accvgpr_write_b32 a7, v1
+; GFX908-NEXT:    v_accvgpr_write_b32 a8, v2
+; GFX908-NEXT:    v_accvgpr_write_b32 a9, v17
+; GFX908-NEXT:    v_mov_b32_e32 v1, s10
+; GFX908-NEXT:    v_mov_b32_e32 v2, s11
+; GFX908-NEXT:    v_mov_b32_e32 v17, s12
+; GFX908-NEXT:    v_accvgpr_write_b32 a10, v1
+; GFX908-NEXT:    v_accvgpr_write_b32 a11, v2
+; GFX908-NEXT:    v_accvgpr_write_b32 a12, v17
+; GFX908-NEXT:    v_mov_b32_e32 v1, s13
+; GFX908-NEXT:    v_mov_b32_e32 v2, s14
+; GFX908-NEXT:    v_mov_b32_e32 v17, s15
+; GFX908-NEXT:    v_accvgpr_write_b32 a13, v1
+; GFX908-NEXT:    v_accvgpr_write_b32 a14, v2
+; GFX908-NEXT:    v_accvgpr_write_b32 a15, v17
+; GFX908-NEXT:    v_mov_b32_e32 v1, 2
+; GFX908-NEXT:    s_nop 1
+; GFX908-NEXT:    v_mfma_i32_32x32x8i8 a[0:15], v0, v1, a[0:15] cbsz:1 abid:2 blgp:3
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 1
+; GFX908-NEXT:    v_accvgpr_read_b32 v15, a15
+; GFX908-NEXT:    v_accvgpr_read_b32 v14, a14
+; GFX908-NEXT:    v_accvgpr_read_b32 v13, a13
+; GFX908-NEXT:    v_accvgpr_read_b32 v12, a12
+; GFX908-NEXT:    v_accvgpr_read_b32 v3, a3
+; GFX908-NEXT:    v_accvgpr_read_b32 v2, a2
+; GFX908-NEXT:    v_accvgpr_read_b32 v1, a1
+; GFX908-NEXT:    v_accvgpr_read_b32 v0, a0
+; GFX908-NEXT:    v_accvgpr_read_b32 v7, a7
+; GFX908-NEXT:    v_accvgpr_read_b32 v6, a6
+; GFX908-NEXT:    v_accvgpr_read_b32 v5, a5
+; GFX908-NEXT:    v_accvgpr_read_b32 v4, a4
+; GFX908-NEXT:    v_accvgpr_read_b32 v11, a11
+; GFX908-NEXT:    v_accvgpr_read_b32 v10, a10
+; GFX908-NEXT:    v_accvgpr_read_b32 v9, a9
+; GFX908-NEXT:    v_accvgpr_read_b32 v8, a8
+; GFX908-NEXT:    global_store_dwordx4 v16, v[12:15], s[16:17] offset:48
+; GFX908-NEXT:    s_nop 0
+; GFX908-NEXT:    global_store_dwordx4 v16, v[8:11], s[16:17] offset:32
+; GFX908-NEXT:    global_store_dwordx4 v16, v[4:7], s[16:17] offset:16
+; GFX908-NEXT:    global_store_dwordx4 v16, v[0:3], s[16:17]
+; GFX908-NEXT:    s_endpgm
+;
+; GFX90A-LABEL: test_mfma_i32_32x32x8i8:
+; GFX90A:       ; %bb.0: ; %bb
+; GFX90A-NEXT:    s_load_dwordx2 s[16:17], s[4:5], 0x24
+; GFX90A-NEXT:    v_mov_b32_e32 v0, 1
+; GFX90A-NEXT:    v_mov_b32_e32 v1, 2
+; GFX90A-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-NEXT:    s_load_dwordx16 s[0:15], s[16:17], 0x0
+; GFX90A-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-NEXT:    v_accvgpr_write_b32 a0, s0
+; GFX90A-NEXT:    v_accvgpr_write_b32 a1, s1
+; GFX90A-NEXT:    v_accvgpr_write_b32 a2, s2
+; GFX90A-NEXT:    v_accvgpr_write_b32 a3, s3
+; GFX90A-NEXT:    v_accvgpr_write_b32 a4, s4
+; GFX90A-NEXT:    v_accvgpr_write_b32 a5, s5
+; GFX90A-NEXT:    v_accvgpr_write_b32 a6, s6
+; GFX90A-NEXT:    v_accvgpr_write_b32 a7, s7
+; GFX90A-NEXT:    v_accvgpr_write_b32 a8, s8
+; GFX90A-NEXT:    v_accvgpr_write_b32 a9, s9
+; GFX90A-NEXT:    v_accvgpr_write_b32 a10, s10
+; GFX90A-NEXT:    v_accvgpr_write_b32 a11, s11
+; GFX90A-NEXT:    v_accvgpr_write_b32 a12, s12
+; GFX90A-NEXT:    v_accvgpr_write_b32 a13, s13
+; GFX90A-NEXT:    v_accvgpr_write_b32 a14, s14
+; GFX90A-NEXT:    v_accvgpr_write_b32 a15, s15
+; GFX90A-NEXT:    s_nop 1
+; GFX90A-NEXT:    v_mfma_i32_32x32x8i8 a[0:15], v0, v1, a[0:15] cbsz:1 abid:2 blgp:3
+; GFX90A-NEXT:    v_mov_b32_e32 v0, 0
+; GFX90A-NEXT:    s_nop 7
+; GFX90A-NEXT:    s_nop 7
+; GFX90A-NEXT:    s_nop 1
+; GFX90A-NEXT:    global_store_dwordx4 v0, a[12:15], s[16:17] offset:48
+; GFX90A-NEXT:    global_store_dwordx4 v0, a[8:11], s[16:17] offset:32
+; GFX90A-NEXT:    global_store_dwordx4 v0, a[4:7], s[16:17] offset:16
+; GFX90A-NEXT:    global_store_dwordx4 v0, a[0:3], s[16:17]
+; GFX90A-NEXT:    s_endpgm
 bb:
   %in.1 = load <16 x i32>, ptr addrspace(1) %arg
   %mai.1 = tail call <16 x i32> @llvm.amdgcn.mfma.i32.32x32x8i8(i32 1, i32 2, <16 x i32> %in.1, i32 1, i32 2, i32 3)
   store <16 x i32> %mai.1, ptr addrspace(1) %arg
   ret void
 }
 
-; GCN-LABEL: {{^}}test_mfma_i32_16x16x16i8:
-; GCN-DAG:        v_mov_b32_e32 [[TWO:v[0-9]+]], 2
-; GCN-DAG:        v_mov_b32_e32 [[ONE:v[0-9]+]], 1
-; GCN:            s_load_dwordx4
-; GFX908-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
-; GFX90A-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
-; GCN:            v_mfma_i32_16x16x16i8 [[RES:a\[[0-9]+:[0-9]+\]]], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
-; GFX908-COUNT-4: v_accvgpr_read_b32
-; GFX908:         global_store_dwordx4
-; GFX90A-NOT:     v_accvgpr_read_b32
-; GFX90A:         global_store_dwordx4 v{{[0-9]+}}, [[RES]]
 define amdgpu_kernel void @test_mfma_i32_16x16x16i8(ptr addrspace(1) %arg) #0 {
+; GFX908-LABEL: test_mfma_i32_16x16x16i8:
+; GFX908:       ; %bb.0: ; %bb
+; GFX908-NEXT:    s_load_dwordx2 s[6:7], s[4:5], 0x24
+; GFX908-NEXT:    v_mov_b32_e32 v0, 1
+; GFX908-NEXT:    v_mov_b32_e32 v1, 2
+; GFX908-NEXT:    v_mov_b32_e32 v4, 0
+; GFX908-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX908-NEXT:    s_load_dwordx4 s[0:3], s[6:7], 0x0
+; GFX908-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX908-NEXT:    v_mov_b32_e32 v5, s0
+; GFX908-NEXT:    v_mov_b32_e32 v2, s1
+; GFX908-NEXT:    v_mov_b32_e32 v3, s2
+; GFX908-NEXT:    v_accvgpr_write_b32 a0, v5
+; GFX908-NEXT:    v_mov_b32_e32 v5, s3
+; GFX908-NEXT:    v_accvgpr_write_b32 a1, v2
+; GFX908-NEXT:    v_accvgpr_write_b32 a2, v3
+; GFX908-NEXT:    v_accvgpr_write_b32 a3, v5
+; GFX908-NEXT:    s_nop 0
+; GFX908-NEXT:    v_mfma_i32_16x16x16i8 a[0:3], v0, v1, a[0:3] cbsz:1 abid:2 blgp:3
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 1
+; GFX908-NEXT:    v_accvgpr_read_b32 v0, a0
+; GFX908-NEXT:    v_accvgpr_read_b32 v1, a1
+; GFX908-NEXT:    v_accvgpr_read_b32 v2, a2
+; GFX908-NEXT:    v_accvgpr_read_b32 v3, a3
+; GFX908-NEXT:    s_nop 1
+; GFX908-NEXT:    global_store_dwordx4 v4, v[0:3], s[6:7]
+; GFX908-NEXT:    s_endpgm
+;
+; GFX90A-LABEL: test_mfma_i32_16x16x16i8:
+; GFX90A:       ; %bb.0: ; %bb
+; GFX90A-NEXT:    s_load_dwordx2 s[6:7], s[4:5], 0x24
+; GFX90A-NEXT:    v_mov_b32_e32 v0, 1
+; GFX90A-NEXT:    v_mov_b32_e32 v2, 2
+; GFX90A-NEXT:    v_mov_b32_e32 v1, 0
+; GFX90A-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-NEXT:    s_load_dwordx4 s[0:3], s[6:7], 0x0
+; GFX90A-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-NEXT:    v_accvgpr_write_b32 a0, s0
+; GFX90A-NEXT:    v_accvgpr_write_b32 a1, s1
+; GFX90A-NEXT:    v_accvgpr_write_b32 a2, s2
+; GFX90A-NEXT:    v_accvgpr_write_b32 a3, s3
+; GFX90A-NEXT:    s_nop 1
+; GFX90A-NEXT:    v_mfma_i32_16x16x16i8 a[0:3], v0, v2, a[0:3] cbsz:1 abid:2 blgp:3
+; GFX90A-NEXT:    s_nop 7
+; GFX90A-NEXT:    s_nop 2
+; GFX90A-NEXT:    global_store_dwordx4 v1, a[0:3], s[6:7]
+; GFX90A-NEXT:    s_endpgm
 bb:
   %in.1 = load <4 x i32>, ptr addrspace(1) %arg
   %mai.1 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.16x16x16i8(i32 1, i32 2, <4 x i32> %in.1, i32 1, i32 2, i32 3)
@@ -44,3 +175,5 @@ bb:
 }
 
 attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }
+;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
+; GCN: {{.*}}

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.xf32.gfx942.ll

@@ -1,40 +1,77 @@
-; RUN: llc -mtriple=amdgcn -mcpu=gfx942 < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX942 %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -global-isel < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GISEL %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -stress-regalloc=10 < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX942 %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -stress-regalloc=10 -global-isel < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GISEL %s
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc -mtriple=amdgcn -mcpu=gfx942 < %s | FileCheck --check-prefixes=GCN,GFX942 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -global-isel < %s | FileCheck --check-prefixes=GCN,GISEL %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -stress-regalloc=10 < %s | FileCheck --check-prefixes=GCN,GFX942 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -stress-regalloc=10 -global-isel < %s | FileCheck --check-prefixes=GCN,GISEL %s
 
 declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x8.xf32(<2 x float>, <2 x float>, <4 x float>, i32, i32, i32)
 declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x4.xf32(<2 x float>, <2 x float>, <16 x float>, i32, i32, i32)
 
-; GCN-LABEL: {{^}}test_mfma_f32_16x16x8xf32:
-; GFX942-DAG:  v_mov_b32_e32 v[[ONE:[0-9]+]], 1.0
-; GFX942-DAG:  v_mov_b32_e32 v[[TWO:[0-9]+]], 2.0
-; GFX942-DAG:  v_mov_b32_e32 v[[THREE:[0-9]+]], 0x40400000
-; GFX942-DAG:  v_mov_b32_e32 v[[FOUR:[0-9]+]], 4.0
-; GCN-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
-; GFX942:      v_mfma_f32_16x16x8_xf32 a[{{[0-9]+:[0-9]+}}], v[[[ONE]]:[[TWO]]], v[[[THREE]]:[[FOUR]]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
-; GISEL:       v_mfma_f32_16x16x8_xf32 a[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
-; GCN-NOT:     v_accvgpr_read_b32
-; GCN:         global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}]
 define amdgpu_kernel void @test_mfma_f32_16x16x8xf32(ptr addrspace(1) %arg) #0 {
+; GFX942-LABEL: test_mfma_f32_16x16x8xf32:
+; GFX942:       ; %bb.0: ; %bb
+; GFX942-NEXT:    s_load_dwordx2 s[6:7], s[4:5], 0x24
+; GFX942-NEXT:    v_mov_b32_e32 v0, 1.0
+; GFX942-NEXT:    v_mov_b32_e32 v1, 2.0
+; GFX942-NEXT:    v_mov_b32_e32 v2, 0x40400000
+; GFX942-NEXT:    v_mov_b32_e32 v3, 4.0
+; GFX942-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-NEXT:    s_load_dwordx4 s[0:3], s[6:7], 0x0
+; GFX942-NEXT:    v_mov_b32_e32 v4, 0
+; GFX942-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-NEXT:    v_accvgpr_write_b32 a0, s0
+; GFX942-NEXT:    v_accvgpr_write_b32 a1, s1
+; GFX942-NEXT:    v_accvgpr_write_b32 a2, s2
+; GFX942-NEXT:    v_accvgpr_write_b32 a3, s3
+; GFX942-NEXT:    s_nop 1
+; GFX942-NEXT:    v_mfma_f32_16x16x8_xf32 a[0:3], v[0:1], v[2:3], a[0:3] cbsz:1 abid:2 blgp:3
+; GFX942-NEXT:    s_nop 6
+; GFX942-NEXT:    global_store_dwordx4 v4, a[0:3], s[6:7]
+; GFX942-NEXT:    s_endpgm
 bb:
   %in.1 = load <4 x float>, ptr addrspace(1) %arg
   %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x8.xf32(<2 x float> <float 1.0, float 2.0>, <2 x float> <float 3.0, float 4.0>, <4 x float> %in.1, i32 1, i32 2, i32 3)
   store <4 x float> %mai.1, ptr addrspace(1) %arg
   ret void
 }
 
-; GCN-LABEL: {{^}}test_mfma_f32_32x32x4xf32:
-; GFX942-DAG:  v_mov_b32_e32 v[[ONE:[0-9]+]], 1.0
-; GFX942-DAG:  v_mov_b32_e32 v[[TWO:[0-9]+]], 2.0
-; GFX942-DAG:  v_mov_b32_e32 v[[THREE:[0-9]+]], 0x40400000
-; GFX942-DAG:  v_mov_b32_e32 v[[FOUR:[0-9]+]], 4.0
-; GCN-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, s{{[0-9]+}}
-; GFX942:      v_mfma_f32_32x32x4_xf32 a[{{[0-9]+:[0-9]+}}], v[[[ONE]]:[[TWO]]], v[[[THREE]]:[[FOUR]]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
-; GISEL:       v_mfma_f32_32x32x4_xf32 a[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}], v[{{[0-9]+:[0-9]+}}], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
-; GCN-NOT:     v_accvgpr_read_b32
-; GCN:         global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}]
 define amdgpu_kernel void @test_mfma_f32_32x32x4xf32(ptr addrspace(1) %arg) #0 {
+; GFX942-LABEL: test_mfma_f32_32x32x4xf32:
+; GFX942:       ; %bb.0: ; %bb
+; GFX942-NEXT:    s_load_dwordx2 s[16:17], s[4:5], 0x24
+; GFX942-NEXT:    v_mov_b32_e32 v0, 1.0
+; GFX942-NEXT:    v_mov_b32_e32 v1, 2.0
+; GFX942-NEXT:    v_mov_b32_e32 v2, 0x40400000
+; GFX942-NEXT:    v_mov_b32_e32 v3, 4.0
+; GFX942-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-NEXT:    s_load_dwordx16 s[0:15], s[16:17], 0x0
+; GFX942-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-NEXT:    v_accvgpr_write_b32 a0, s0
+; GFX942-NEXT:    v_accvgpr_write_b32 a1, s1
+; GFX942-NEXT:    v_accvgpr_write_b32 a2, s2
+; GFX942-NEXT:    v_accvgpr_write_b32 a3, s3
+; GFX942-NEXT:    v_accvgpr_write_b32 a4, s4
+; GFX942-NEXT:    v_accvgpr_write_b32 a5, s5
+; GFX942-NEXT:    v_accvgpr_write_b32 a6, s6
+; GFX942-NEXT:    v_accvgpr_write_b32 a7, s7
+; GFX942-NEXT:    v_accvgpr_write_b32 a8, s8
+; GFX942-NEXT:    v_accvgpr_write_b32 a9, s9
+; GFX942-NEXT:    v_accvgpr_write_b32 a10, s10
+; GFX942-NEXT:    v_accvgpr_write_b32 a11, s11
+; GFX942-NEXT:    v_accvgpr_write_b32 a12, s12
+; GFX942-NEXT:    v_accvgpr_write_b32 a13, s13
+; GFX942-NEXT:    v_accvgpr_write_b32 a14, s14
+; GFX942-NEXT:    v_accvgpr_write_b32 a15, s15
+; GFX942-NEXT:    s_nop 1
+; GFX942-NEXT:    v_mfma_f32_32x32x4_xf32 a[0:15], v[0:1], v[2:3], a[0:15] cbsz:1 abid:2 blgp:3
+; GFX942-NEXT:    v_mov_b32_e32 v0, 0
+; GFX942-NEXT:    s_nop 7
+; GFX942-NEXT:    s_nop 1
+; GFX942-NEXT:    global_store_dwordx4 v0, a[12:15], s[16:17] offset:48
+; GFX942-NEXT:    global_store_dwordx4 v0, a[8:11], s[16:17] offset:32
+; GFX942-NEXT:    global_store_dwordx4 v0, a[4:7], s[16:17] offset:16
+; GFX942-NEXT:    global_store_dwordx4 v0, a[0:3], s[16:17]
+; GFX942-NEXT:    s_endpgm
 bb:
   %in.1 = load <16 x float>, ptr addrspace(1) %arg
   %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x4.xf32(<2 x float> <float 1.0, float 2.0>, <2 x float> <float 3.0, float 4.0>, <16 x float> %in.1, i32 1, i32 2, i32 3)
@@ -43,3 +80,6 @@ bb:
 }
 
 attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }
+;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
+; GCN: {{.*}}
+; GISEL: {{.*}}