Skip to content

Commit 4dc5ca9

Browse files
arsenmgithub-actions[bot]
arsenm
authored and
github-actions[bot]
committed
Automerge: AMDGPU: Fix test with broken checks
This was broken in 6118a25 update_llc_test_checks behavior in the conflict case is dangerous and terrible; it silently deletes all checks and inserts the unused check prefixes at the bottom.
2 parents eb08ecd + 1c6e75c commit 4dc5ca9

File tree

1 file changed

+243
-6
lines changed

1 file changed

+243
-6
lines changed

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

Lines changed: 243 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1,13 +1,100 @@
11
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2-
; RUN: llc -mtriple=amdgcn -mcpu=gfx942 < %s | FileCheck --check-prefixes=GCN,GFX942 %s
3-
; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -global-isel < %s | FileCheck --check-prefixes=GCN,GISEL %s
4-
; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -stress-regalloc=10 < %s | FileCheck --check-prefixes=GCN,GFX942 %s
5-
; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -stress-regalloc=10 -global-isel < %s | FileCheck --check-prefixes=GCN,GISEL %s
2+
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx942 < %s | FileCheck --check-prefixes=GFX942,GFX942-SDAG %s
3+
; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx942 < %s | FileCheck --check-prefixes=GFX942,GFX942-GISEL %s
4+
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx942 -stress-regalloc=10 < %s | FileCheck --check-prefixes=GFX942-STRESS,GFX942-SDAG-STRESS %s
5+
; RUN: llc -global-isel -mtriple=amdgcn -mcpu=gfx942 -stress-regalloc=10 < %s | FileCheck --check-prefixes=GFX942-STRESS,GFX942-GISEL-STRESS %s
66

77
declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x8.xf32(<2 x float>, <2 x float>, <4 x float>, i32, i32, i32)
88
declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x4.xf32(<2 x float>, <2 x float>, <16 x float>, i32, i32, i32)
99

1010
define amdgpu_kernel void @test_mfma_f32_16x16x8xf32(ptr addrspace(1) %arg) #0 {
11+
; GFX942-SDAG-LABEL: test_mfma_f32_16x16x8xf32:
12+
; GFX942-SDAG: ; %bb.0: ; %bb
13+
; GFX942-SDAG-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
14+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v4, 1.0
15+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v5, 2.0
16+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v0, 0x40400000
17+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v1, 4.0
18+
; GFX942-SDAG-NEXT: s_waitcnt lgkmcnt(0)
19+
; GFX942-SDAG-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0
20+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v2, 0
21+
; GFX942-SDAG-NEXT: s_waitcnt lgkmcnt(0)
22+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a0, s0
23+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a1, s1
24+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a2, s2
25+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a3, s3
26+
; GFX942-SDAG-NEXT: s_nop 1
27+
; GFX942-SDAG-NEXT: v_mfma_f32_16x16x8_xf32 a[0:3], v[4:5], v[0:1], a[0:3] cbsz:1 abid:2 blgp:3
28+
; GFX942-SDAG-NEXT: s_nop 6
29+
; GFX942-SDAG-NEXT: global_store_dwordx4 v2, a[0:3], s[6:7]
30+
; GFX942-SDAG-NEXT: s_endpgm
31+
;
32+
; GFX942-GISEL-LABEL: test_mfma_f32_16x16x8xf32:
33+
; GFX942-GISEL: ; %bb.0: ; %bb
34+
; GFX942-GISEL-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
35+
; GFX942-GISEL-NEXT: s_mov_b32 s4, 1.0
36+
; GFX942-GISEL-NEXT: s_mov_b32 s5, 2.0
37+
; GFX942-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[4:5]
38+
; GFX942-GISEL-NEXT: s_mov_b32 s4, 0x40400000
39+
; GFX942-GISEL-NEXT: s_waitcnt lgkmcnt(0)
40+
; GFX942-GISEL-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0
41+
; GFX942-GISEL-NEXT: s_mov_b32 s5, 4.0
42+
; GFX942-GISEL-NEXT: v_mov_b64_e32 v[2:3], s[4:5]
43+
; GFX942-GISEL-NEXT: s_waitcnt lgkmcnt(0)
44+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a0, s0
45+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a1, s1
46+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a2, s2
47+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a3, s3
48+
; GFX942-GISEL-NEXT: s_nop 1
49+
; GFX942-GISEL-NEXT: v_mfma_f32_16x16x8_xf32 a[0:3], v[0:1], v[2:3], a[0:3] cbsz:1 abid:2 blgp:3
50+
; GFX942-GISEL-NEXT: v_mov_b32_e32 v0, 0
51+
; GFX942-GISEL-NEXT: s_nop 5
52+
; GFX942-GISEL-NEXT: global_store_dwordx4 v0, a[0:3], s[6:7]
53+
; GFX942-GISEL-NEXT: s_endpgm
54+
;
55+
; GFX942-SDAG-STRESS-LABEL: test_mfma_f32_16x16x8xf32:
56+
; GFX942-SDAG-STRESS: ; %bb.0: ; %bb
57+
; GFX942-SDAG-STRESS-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
58+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v0, 1.0
59+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v1, 2.0
60+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v2, 0x40400000
61+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v3, 4.0
62+
; GFX942-SDAG-STRESS-NEXT: s_waitcnt lgkmcnt(0)
63+
; GFX942-SDAG-STRESS-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0
64+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v4, 0
65+
; GFX942-SDAG-STRESS-NEXT: s_waitcnt lgkmcnt(0)
66+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a0, s0
67+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a1, s1
68+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a2, s2
69+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a3, s3
70+
; GFX942-SDAG-STRESS-NEXT: s_nop 1
71+
; GFX942-SDAG-STRESS-NEXT: v_mfma_f32_16x16x8_xf32 a[0:3], v[0:1], v[2:3], a[0:3] cbsz:1 abid:2 blgp:3
72+
; GFX942-SDAG-STRESS-NEXT: s_nop 6
73+
; GFX942-SDAG-STRESS-NEXT: global_store_dwordx4 v4, a[0:3], s[6:7]
74+
; GFX942-SDAG-STRESS-NEXT: s_endpgm
75+
;
76+
; GFX942-GISEL-STRESS-LABEL: test_mfma_f32_16x16x8xf32:
77+
; GFX942-GISEL-STRESS: ; %bb.0: ; %bb
78+
; GFX942-GISEL-STRESS-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
79+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s0, 1.0
80+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s2, 0x40400000
81+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s1, 2.0
82+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s3, 4.0
83+
; GFX942-GISEL-STRESS-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
84+
; GFX942-GISEL-STRESS-NEXT: v_mov_b64_e32 v[2:3], s[2:3]
85+
; GFX942-GISEL-STRESS-NEXT: s_waitcnt lgkmcnt(0)
86+
; GFX942-GISEL-STRESS-NEXT: s_load_dwordx4 s[0:3], s[6:7], 0x0
87+
; GFX942-GISEL-STRESS-NEXT: s_waitcnt lgkmcnt(0)
88+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a0, s0
89+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a1, s1
90+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a2, s2
91+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a3, s3
92+
; GFX942-GISEL-STRESS-NEXT: s_nop 1
93+
; GFX942-GISEL-STRESS-NEXT: v_mfma_f32_16x16x8_xf32 a[0:3], v[0:1], v[2:3], a[0:3] cbsz:1 abid:2 blgp:3
94+
; GFX942-GISEL-STRESS-NEXT: v_mov_b32_e32 v0, 0
95+
; GFX942-GISEL-STRESS-NEXT: s_nop 5
96+
; GFX942-GISEL-STRESS-NEXT: global_store_dwordx4 v0, a[0:3], s[6:7]
97+
; GFX942-GISEL-STRESS-NEXT: s_endpgm
1198
bb:
1299
%in.1 = load <4 x float>, ptr addrspace(1) %arg
13100
%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)
@@ -16,6 +103,157 @@ bb:
16103
}
17104

18105
define amdgpu_kernel void @test_mfma_f32_32x32x4xf32(ptr addrspace(1) %arg) #0 {
106+
; GFX942-SDAG-LABEL: test_mfma_f32_32x32x4xf32:
107+
; GFX942-SDAG: ; %bb.0: ; %bb
108+
; GFX942-SDAG-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24
109+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v2, 1.0
110+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v3, 2.0
111+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v0, 0x40400000
112+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v1, 4.0
113+
; GFX942-SDAG-NEXT: s_waitcnt lgkmcnt(0)
114+
; GFX942-SDAG-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0
115+
; GFX942-SDAG-NEXT: s_waitcnt lgkmcnt(0)
116+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a0, s0
117+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a1, s1
118+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a2, s2
119+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a3, s3
120+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a4, s4
121+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a5, s5
122+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a6, s6
123+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a7, s7
124+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a8, s8
125+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a9, s9
126+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a10, s10
127+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a11, s11
128+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a12, s12
129+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a13, s13
130+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a14, s14
131+
; GFX942-SDAG-NEXT: v_accvgpr_write_b32 a15, s15
132+
; GFX942-SDAG-NEXT: s_nop 1
133+
; GFX942-SDAG-NEXT: v_mfma_f32_32x32x4_xf32 a[0:15], v[2:3], v[0:1], a[0:15] cbsz:1 abid:2 blgp:3
134+
; GFX942-SDAG-NEXT: v_mov_b32_e32 v0, 0
135+
; GFX942-SDAG-NEXT: s_nop 7
136+
; GFX942-SDAG-NEXT: s_nop 1
137+
; GFX942-SDAG-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48
138+
; GFX942-SDAG-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32
139+
; GFX942-SDAG-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16
140+
; GFX942-SDAG-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]
141+
; GFX942-SDAG-NEXT: s_endpgm
142+
;
143+
; GFX942-GISEL-LABEL: test_mfma_f32_32x32x4xf32:
144+
; GFX942-GISEL: ; %bb.0: ; %bb
145+
; GFX942-GISEL-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24
146+
; GFX942-GISEL-NEXT: s_mov_b32 s18, 1.0
147+
; GFX942-GISEL-NEXT: s_mov_b32 s19, 2.0
148+
; GFX942-GISEL-NEXT: v_mov_b64_e32 v[0:1], s[18:19]
149+
; GFX942-GISEL-NEXT: s_mov_b32 s18, 0x40400000
150+
; GFX942-GISEL-NEXT: s_waitcnt lgkmcnt(0)
151+
; GFX942-GISEL-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0
152+
; GFX942-GISEL-NEXT: s_mov_b32 s19, 4.0
153+
; GFX942-GISEL-NEXT: v_mov_b64_e32 v[2:3], s[18:19]
154+
; GFX942-GISEL-NEXT: s_waitcnt lgkmcnt(0)
155+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a0, s0
156+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a1, s1
157+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a2, s2
158+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a3, s3
159+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a4, s4
160+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a5, s5
161+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a6, s6
162+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a7, s7
163+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a8, s8
164+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a9, s9
165+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a10, s10
166+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a11, s11
167+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a12, s12
168+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a13, s13
169+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a14, s14
170+
; GFX942-GISEL-NEXT: v_accvgpr_write_b32 a15, s15
171+
; GFX942-GISEL-NEXT: s_nop 1
172+
; GFX942-GISEL-NEXT: v_mfma_f32_32x32x4_xf32 a[0:15], v[0:1], v[2:3], a[0:15] cbsz:1 abid:2 blgp:3
173+
; GFX942-GISEL-NEXT: v_mov_b32_e32 v0, 0
174+
; GFX942-GISEL-NEXT: s_nop 7
175+
; GFX942-GISEL-NEXT: s_nop 1
176+
; GFX942-GISEL-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]
177+
; GFX942-GISEL-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16
178+
; GFX942-GISEL-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32
179+
; GFX942-GISEL-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48
180+
; GFX942-GISEL-NEXT: s_endpgm
181+
;
182+
; GFX942-SDAG-STRESS-LABEL: test_mfma_f32_32x32x4xf32:
183+
; GFX942-SDAG-STRESS: ; %bb.0: ; %bb
184+
; GFX942-SDAG-STRESS-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24
185+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v0, 1.0
186+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v1, 2.0
187+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v2, 0x40400000
188+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v3, 4.0
189+
; GFX942-SDAG-STRESS-NEXT: s_waitcnt lgkmcnt(0)
190+
; GFX942-SDAG-STRESS-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0
191+
; GFX942-SDAG-STRESS-NEXT: s_waitcnt lgkmcnt(0)
192+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a0, s0
193+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a1, s1
194+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a2, s2
195+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a3, s3
196+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a4, s4
197+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a5, s5
198+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a6, s6
199+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a7, s7
200+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a8, s8
201+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a9, s9
202+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a10, s10
203+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a11, s11
204+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a12, s12
205+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a13, s13
206+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a14, s14
207+
; GFX942-SDAG-STRESS-NEXT: v_accvgpr_write_b32 a15, s15
208+
; GFX942-SDAG-STRESS-NEXT: s_nop 1
209+
; GFX942-SDAG-STRESS-NEXT: v_mfma_f32_32x32x4_xf32 a[0:15], v[0:1], v[2:3], a[0:15] cbsz:1 abid:2 blgp:3
210+
; GFX942-SDAG-STRESS-NEXT: v_mov_b32_e32 v0, 0
211+
; GFX942-SDAG-STRESS-NEXT: s_nop 7
212+
; GFX942-SDAG-STRESS-NEXT: s_nop 1
213+
; GFX942-SDAG-STRESS-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48
214+
; GFX942-SDAG-STRESS-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32
215+
; GFX942-SDAG-STRESS-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16
216+
; GFX942-SDAG-STRESS-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]
217+
; GFX942-SDAG-STRESS-NEXT: s_endpgm
218+
;
219+
; GFX942-GISEL-STRESS-LABEL: test_mfma_f32_32x32x4xf32:
220+
; GFX942-GISEL-STRESS: ; %bb.0: ; %bb
221+
; GFX942-GISEL-STRESS-NEXT: s_load_dwordx2 s[16:17], s[4:5], 0x24
222+
; GFX942-GISEL-STRESS-NEXT: s_waitcnt lgkmcnt(0)
223+
; GFX942-GISEL-STRESS-NEXT: s_load_dwordx16 s[0:15], s[16:17], 0x0
224+
; GFX942-GISEL-STRESS-NEXT: s_waitcnt lgkmcnt(0)
225+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a0, s0
226+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a1, s1
227+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a2, s2
228+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a3, s3
229+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a4, s4
230+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a5, s5
231+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a6, s6
232+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a7, s7
233+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a8, s8
234+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a9, s9
235+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a10, s10
236+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a11, s11
237+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a12, s12
238+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a13, s13
239+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a14, s14
240+
; GFX942-GISEL-STRESS-NEXT: v_accvgpr_write_b32 a15, s15
241+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s0, 1.0
242+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s1, 2.0
243+
; GFX942-GISEL-STRESS-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
244+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s0, 0x40400000
245+
; GFX942-GISEL-STRESS-NEXT: s_mov_b32 s1, 4.0
246+
; GFX942-GISEL-STRESS-NEXT: v_mov_b64_e32 v[2:3], s[0:1]
247+
; GFX942-GISEL-STRESS-NEXT: s_nop 1
248+
; GFX942-GISEL-STRESS-NEXT: v_mfma_f32_32x32x4_xf32 a[0:15], v[0:1], v[2:3], a[0:15] cbsz:1 abid:2 blgp:3
249+
; GFX942-GISEL-STRESS-NEXT: v_mov_b32_e32 v0, 0
250+
; GFX942-GISEL-STRESS-NEXT: s_nop 7
251+
; GFX942-GISEL-STRESS-NEXT: s_nop 1
252+
; GFX942-GISEL-STRESS-NEXT: global_store_dwordx4 v0, a[0:3], s[16:17]
253+
; GFX942-GISEL-STRESS-NEXT: global_store_dwordx4 v0, a[4:7], s[16:17] offset:16
254+
; GFX942-GISEL-STRESS-NEXT: global_store_dwordx4 v0, a[8:11], s[16:17] offset:32
255+
; GFX942-GISEL-STRESS-NEXT: global_store_dwordx4 v0, a[12:15], s[16:17] offset:48
256+
; GFX942-GISEL-STRESS-NEXT: s_endpgm
19257
bb:
20258
%in.1 = load <16 x float>, ptr addrspace(1) %arg
21259
%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)
@@ -25,6 +263,5 @@ bb:
25263

26264
attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }
27265
;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
28-
; GCN: {{.*}}
29266
; GFX942: {{.*}}
30-
; GISEL: {{.*}}
267+
; GFX942-STRESS: {{.*}}

0 commit comments

Comments
 (0)