CodeGen/AMDGPU/llvm.amdgcn.mfma.ll

1 ; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
25 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32:
26 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
27 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
28 ; GCN-DAG: s_load_dwordx16
29 ; GCN-DAG: s_load_dwordx16
30 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
31 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
32 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
33 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
34 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
35 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
36 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
37 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
38 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
39 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
40 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
41 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
42 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
43 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
44 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
45 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
46 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
47 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
48 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
49 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
50 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
51 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
52 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
53 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
54 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
55 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
56 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
57 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
58 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
59 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
60 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
61 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
62 ; GCN: v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 ab…
63 ; GCN-DAG: v_accvgpr_read_b32
64 ; GCN-DAG: v_accvgpr_read_b32
65 ; GCN-DAG: v_accvgpr_read_b32
66 ; GCN-DAG: v_accvgpr_read_b32
67 ; GCN-DAG: v_accvgpr_read_b32
68 ; GCN-DAG: v_accvgpr_read_b32
69 ; GCN-DAG: v_accvgpr_read_b32
70 ; GCN-DAG: v_accvgpr_read_b32
71 ; GCN-DAG: v_accvgpr_read_b32
72 ; GCN-DAG: v_accvgpr_read_b32
73 ; GCN-DAG: v_accvgpr_read_b32
74 ; GCN-DAG: v_accvgpr_read_b32
75 ; GCN-DAG: v_accvgpr_read_b32
76 ; GCN-DAG: v_accvgpr_read_b32
77 ; GCN-DAG: v_accvgpr_read_b32
78 ; GCN-DAG: v_accvgpr_read_b32
79 ; GCN-DAG: v_accvgpr_read_b32
80 ; GCN-DAG: v_accvgpr_read_b32
81 ; GCN-DAG: v_accvgpr_read_b32
82 ; GCN-DAG: v_accvgpr_read_b32
83 ; GCN-DAG: v_accvgpr_read_b32
84 ; GCN-DAG: v_accvgpr_read_b32
85 ; GCN-DAG: v_accvgpr_read_b32
86 ; GCN-DAG: v_accvgpr_read_b32
87 ; GCN-DAG: v_accvgpr_read_b32
88 ; GCN-DAG: v_accvgpr_read_b32
89 ; GCN-DAG: v_accvgpr_read_b32
90 ; GCN-DAG: v_accvgpr_read_b32
91 ; GCN-DAG: v_accvgpr_read_b32
92 ; GCN-DAG: v_accvgpr_read_b32
93 ; GCN-DAG: v_accvgpr_read_b32
94 ; GCN-DAG: v_accvgpr_read_b32
95 ; GCN-DAG: global_store_dwordx4
96 ; GCN-DAG: global_store_dwordx4
97 ; GCN-DAG: global_store_dwordx4
98 ; GCN-DAG: global_store_dwordx4
99 ; GCN-DAG: global_store_dwordx4
100 ; GCN-DAG: global_store_dwordx4
101 ; GCN-DAG: global_store_dwordx4
102 ; GCN-DAG: global_store_dwordx4
111 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x1f32:
112 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
113 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
114 ; GCN: s_load_dwordx16
115 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
116 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
117 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
118 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
119 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
120 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
121 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
122 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
123 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
124 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
125 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
126 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
127 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
128 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
129 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
130 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
131 ; GCN: v_mfma_f32_16x16x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 ab…
132 ; GCN-DAG: v_accvgpr_read_b32
133 ; GCN-DAG: v_accvgpr_read_b32
134 ; GCN-DAG: v_accvgpr_read_b32
135 ; GCN-DAG: v_accvgpr_read_b32
136 ; GCN-DAG: v_accvgpr_read_b32
137 ; GCN-DAG: v_accvgpr_read_b32
138 ; GCN-DAG: v_accvgpr_read_b32
139 ; GCN-DAG: v_accvgpr_read_b32
140 ; GCN-DAG: v_accvgpr_read_b32
141 ; GCN-DAG: v_accvgpr_read_b32
142 ; GCN-DAG: v_accvgpr_read_b32
143 ; GCN-DAG: v_accvgpr_read_b32
144 ; GCN-DAG: v_accvgpr_read_b32
145 ; GCN-DAG: v_accvgpr_read_b32
146 ; GCN-DAG: v_accvgpr_read_b32
147 ; GCN-DAG: v_accvgpr_read_b32
148 ; GCN-DAG: global_store_dwordx4
149 ; GCN-DAG: global_store_dwordx4
150 ; GCN-DAG: global_store_dwordx4
151 ; GCN-DAG: global_store_dwordx4
160 ; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32:
161 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
162 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
163 ; GCN: s_load_dwordx4
164 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
165 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
166 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
167 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
168 ; GCN: v_mfma_f32_4x4x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid…
169 ; GCN: v_accvgpr_read_b32
170 ; GCN: v_accvgpr_read_b32
171 ; GCN: v_accvgpr_read_b32
172 ; GCN: v_accvgpr_read_b32
173 ; GCN: global_store_dwordx4
182 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x2f32:
183 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
184 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
185 ; GCN: s_load_dwordx16
186 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
187 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
188 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
189 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
190 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
191 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
192 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
193 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
194 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
195 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
196 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
197 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
198 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
199 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
200 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
201 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
202 ; GCN: v_mfma_f32_32x32x2f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 ab…
203 ; GCN-DAG: v_accvgpr_read_b32
204 ; GCN-DAG: v_accvgpr_read_b32
205 ; GCN-DAG: v_accvgpr_read_b32
206 ; GCN-DAG: v_accvgpr_read_b32
207 ; GCN-DAG: v_accvgpr_read_b32
208 ; GCN-DAG: v_accvgpr_read_b32
209 ; GCN-DAG: v_accvgpr_read_b32
210 ; GCN-DAG: v_accvgpr_read_b32
211 ; GCN-DAG: v_accvgpr_read_b32
212 ; GCN-DAG: v_accvgpr_read_b32
213 ; GCN-DAG: v_accvgpr_read_b32
214 ; GCN-DAG: v_accvgpr_read_b32
215 ; GCN-DAG: v_accvgpr_read_b32
216 ; GCN-DAG: v_accvgpr_read_b32
217 ; GCN-DAG: v_accvgpr_read_b32
218 ; GCN-DAG: v_accvgpr_read_b32
219 ; GCN-DAG: global_store_dwordx4
220 ; GCN-DAG: global_store_dwordx4
221 ; GCN-DAG: global_store_dwordx4
222 ; GCN-DAG: global_store_dwordx4
231 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x4f32:
232 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
233 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
234 ; GCN: s_load_dwordx4
235 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
236 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
237 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
238 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
239 ; GCN: v_mfma_f32_16x16x4f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 ab…
240 ; GCN-DAG: v_accvgpr_read_b32
241 ; GCN-DAG: v_accvgpr_read_b32
242 ; GCN-DAG: v_accvgpr_read_b32
243 ; GCN-DAG: v_accvgpr_read_b32
244 ; GCN-DAG: global_store_dwordx4
253 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x4f16:
254 ; GCN-DAG: s_load_dwordx16
255 ; GCN-DAG: s_load_dwordx16
256 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
257 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
258 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
259 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
260 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
261 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
262 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
263 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
264 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
265 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
266 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
267 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
268 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
269 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
270 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
271 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
272 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
273 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
274 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
275 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
276 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
277 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
278 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
279 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
280 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
281 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
282 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
283 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
284 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
285 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
286 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
287 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
288 ; GCN: v_mfma_f32_32x32x4f16 a[{{[0-9]+:[0-9]+}}], {{v\[[0-9]+:[0-9]+\]}}, {{v\[[0-9]+:[0-9]+\]}}, …
289 ; GCN-DAG: v_accvgpr_read_b32
290 ; GCN-DAG: v_accvgpr_read_b32
291 ; GCN-DAG: v_accvgpr_read_b32
292 ; GCN-DAG: v_accvgpr_read_b32
293 ; GCN-DAG: v_accvgpr_read_b32
294 ; GCN-DAG: v_accvgpr_read_b32
295 ; GCN-DAG: v_accvgpr_read_b32
296 ; GCN-DAG: v_accvgpr_read_b32
297 ; GCN-DAG: v_accvgpr_read_b32
298 ; GCN-DAG: v_accvgpr_read_b32
299 ; GCN-DAG: v_accvgpr_read_b32
300 ; GCN-DAG: v_accvgpr_read_b32
301 ; GCN-DAG: v_accvgpr_read_b32
302 ; GCN-DAG: v_accvgpr_read_b32
303 ; GCN-DAG: v_accvgpr_read_b32
304 ; GCN-DAG: v_accvgpr_read_b32
305 ; GCN-DAG: v_accvgpr_read_b32
306 ; GCN-DAG: v_accvgpr_read_b32
307 ; GCN-DAG: v_accvgpr_read_b32
308 ; GCN-DAG: v_accvgpr_read_b32
309 ; GCN-DAG: v_accvgpr_read_b32
310 ; GCN-DAG: v_accvgpr_read_b32
311 ; GCN-DAG: v_accvgpr_read_b32
312 ; GCN-DAG: v_accvgpr_read_b32
313 ; GCN-DAG: v_accvgpr_read_b32
314 ; GCN-DAG: v_accvgpr_read_b32
315 ; GCN-DAG: v_accvgpr_read_b32
316 ; GCN-DAG: v_accvgpr_read_b32
317 ; GCN-DAG: v_accvgpr_read_b32
318 ; GCN-DAG: v_accvgpr_read_b32
319 ; GCN-DAG: v_accvgpr_read_b32
320 ; GCN-DAG: v_accvgpr_read_b32
321 ; GCN-DAG: global_store_dwordx4
322 ; GCN-DAG: global_store_dwordx4
323 ; GCN-DAG: global_store_dwordx4
324 ; GCN-DAG: global_store_dwordx4
325 ; GCN-DAG: global_store_dwordx4
326 ; GCN-DAG: global_store_dwordx4
327 ; GCN-DAG: global_store_dwordx4
328 ; GCN-DAG: global_store_dwordx4
340 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x4f16:
341 ; GCN: s_load_dwordx16
342 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
343 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
344 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
345 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
346 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
347 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
348 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
349 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
350 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
351 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
352 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
353 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
354 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
355 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
356 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
357 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
358 ; GCN: v_mfma_f32_16x16x4f16 a[{{[0-9]+:[0-9]+}}], {{v\[[0-9]+:[0-9]+\]}}, {{v\[[0-9]+:[0-9]+\]}}, …
359 ; GCN-DAG: v_accvgpr_read_b32
360 ; GCN-DAG: v_accvgpr_read_b32
361 ; GCN-DAG: v_accvgpr_read_b32
362 ; GCN-DAG: v_accvgpr_read_b32
363 ; GCN-DAG: v_accvgpr_read_b32
364 ; GCN-DAG: v_accvgpr_read_b32
365 ; GCN-DAG: v_accvgpr_read_b32
366 ; GCN-DAG: v_accvgpr_read_b32
367 ; GCN-DAG: v_accvgpr_read_b32
368 ; GCN-DAG: v_accvgpr_read_b32
369 ; GCN-DAG: v_accvgpr_read_b32
370 ; GCN-DAG: v_accvgpr_read_b32
371 ; GCN-DAG: v_accvgpr_read_b32
372 ; GCN-DAG: v_accvgpr_read_b32
373 ; GCN-DAG: v_accvgpr_read_b32
374 ; GCN-DAG: v_accvgpr_read_b32
375 ; GCN-DAG: global_store_dwordx4
376 ; GCN-DAG: global_store_dwordx4
377 ; GCN-DAG: global_store_dwordx4
378 ; GCN-DAG: global_store_dwordx4
390 ; GCN-LABEL: {{^}}test_mfma_f32_4x4x4f16:
391 ; GCN: s_load_dwordx4
392 ; GCN: s_load_dwordx2
393 ; GCN: s_load_dwordx2
394 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
395 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
396 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
397 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
398 ; GCN: v_mfma_f32_4x4x4f16 a[{{[0-9]+:[0-9]+}}], {{v\[[0-9]+:[0-9]+\]}}, {{v\[[0-9]+:[0-9]+\]}}, a[…
399 ; GCN-DAG: v_accvgpr_read_b32
400 ; GCN-DAG: v_accvgpr_read_b32
401 ; GCN-DAG: v_accvgpr_read_b32
402 ; GCN-DAG: v_accvgpr_read_b32
403 ; GCN-DAG: global_store_dwordx4
415 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x8f16:
416 ; GCN: s_load_dwordx16
417 ; GCN: s_waitcnt lgkmcnt(0)
418 ; GCN: v_mov_b32_e32 v{{[0-9]+}}, s{{[0-9]+}}
419 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
420 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
421 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
422 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
423 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
424 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
425 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
426 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
427 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
428 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
429 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
430 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
431 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
432 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
433 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
434 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
435 ; GCN: v_mfma_f32_32x32x8f16 a[{{[0-9]+:[0-9]+}}], {{v\[[0-9]+:[0-9]+\]}}, {{v\[[0-9]+:[0-9]+\]}}, …
436 ; GCN-DAG: v_accvgpr_read_b32
437 ; GCN-DAG: v_accvgpr_read_b32
438 ; GCN-DAG: v_accvgpr_read_b32
439 ; GCN-DAG: v_accvgpr_read_b32
440 ; GCN-DAG: v_accvgpr_read_b32
441 ; GCN-DAG: v_accvgpr_read_b32
442 ; GCN-DAG: v_accvgpr_read_b32
443 ; GCN-DAG: v_accvgpr_read_b32
444 ; GCN-DAG: v_accvgpr_read_b32
445 ; GCN-DAG: v_accvgpr_read_b32
446 ; GCN-DAG: v_accvgpr_read_b32
447 ; GCN-DAG: v_accvgpr_read_b32
448 ; GCN-DAG: v_accvgpr_read_b32
449 ; GCN-DAG: v_accvgpr_read_b32
450 ; GCN-DAG: v_accvgpr_read_b32
451 ; GCN-DAG: v_accvgpr_read_b32
452 ; GCN-DAG: global_store_dwordx4
453 ; GCN-DAG: global_store_dwordx4
454 ; GCN-DAG: global_store_dwordx4
455 ; GCN-DAG: global_store_dwordx4
467 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x16f16:
468 ; GCN: s_load_dwordx4
469 ; GCN: s_load_dwordx4
470 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
471 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
472 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
473 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
474 ; GCN: v_mfma_f32_16x16x16f16 a[{{[0-9]+:[0-9]+}}], {{v\[[0-9]+:[0-9]+\]}}, {{v\[[0-9]+:[0-9]+\]}},…
475 ; GCN-DAG: v_accvgpr_read_b32
476 ; GCN-DAG: v_accvgpr_read_b32
477 ; GCN-DAG: v_accvgpr_read_b32
478 ; GCN-DAG: v_accvgpr_read_b32
479 ; GCN-DAG: global_store_dwordx4
491 ; GCN-LABEL: {{^}}test_mfma_i32_32x32x4i8:
492 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2
493 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1
494 ; GCN-DAG: s_load_dwordx16
495 ; GCN-DAG: s_load_dwordx16
496 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
497 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
498 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
499 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
500 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
501 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
502 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
503 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
504 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
505 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
506 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
507 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
508 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
509 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
510 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
511 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
512 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
513 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
514 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
515 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
516 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
517 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
518 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
519 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
520 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
521 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
522 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
523 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
524 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
525 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
526 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
527 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
528 ; GCN: v_mfma_i32_32x32x4i8 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abi…
529 ; GCN-DAG: v_accvgpr_read_b32
530 ; GCN-DAG: v_accvgpr_read_b32
531 ; GCN-DAG: v_accvgpr_read_b32
532 ; GCN-DAG: v_accvgpr_read_b32
533 ; GCN-DAG: v_accvgpr_read_b32
534 ; GCN-DAG: v_accvgpr_read_b32
535 ; GCN-DAG: v_accvgpr_read_b32
536 ; GCN-DAG: v_accvgpr_read_b32
537 ; GCN-DAG: v_accvgpr_read_b32
538 ; GCN-DAG: v_accvgpr_read_b32
539 ; GCN-DAG: v_accvgpr_read_b32
540 ; GCN-DAG: v_accvgpr_read_b32
541 ; GCN-DAG: v_accvgpr_read_b32
542 ; GCN-DAG: v_accvgpr_read_b32
543 ; GCN-DAG: v_accvgpr_read_b32
544 ; GCN-DAG: v_accvgpr_read_b32
545 ; GCN-DAG: v_accvgpr_read_b32
546 ; GCN-DAG: v_accvgpr_read_b32
547 ; GCN-DAG: v_accvgpr_read_b32
548 ; GCN-DAG: v_accvgpr_read_b32
549 ; GCN-DAG: v_accvgpr_read_b32
550 ; GCN-DAG: v_accvgpr_read_b32
551 ; GCN-DAG: v_accvgpr_read_b32
552 ; GCN-DAG: v_accvgpr_read_b32
553 ; GCN-DAG: v_accvgpr_read_b32
554 ; GCN-DAG: v_accvgpr_read_b32
555 ; GCN-DAG: v_accvgpr_read_b32
556 ; GCN-DAG: v_accvgpr_read_b32
557 ; GCN-DAG: v_accvgpr_read_b32
558 ; GCN-DAG: v_accvgpr_read_b32
559 ; GCN-DAG: v_accvgpr_read_b32
560 ; GCN-DAG: v_accvgpr_read_b32
561 ; GCN-DAG: global_store_dwordx4
562 ; GCN-DAG: global_store_dwordx4
563 ; GCN-DAG: global_store_dwordx4
564 ; GCN-DAG: global_store_dwordx4
565 ; GCN-DAG: global_store_dwordx4
566 ; GCN-DAG: global_store_dwordx4
567 ; GCN-DAG: global_store_dwordx4
568 ; GCN-DAG: global_store_dwordx4
577 ; GCN-LABEL: {{^}}test_mfma_i32_16x16x4i8:
578 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2
579 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1
580 ; GCN: s_load_dwordx16
581 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
582 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
583 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
584 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
585 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
586 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
587 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
588 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
589 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
590 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
591 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
592 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
593 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
594 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
595 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
596 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
597 ; GCN: v_mfma_i32_16x16x4i8 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abi…
598 ; GCN-DAG: v_accvgpr_read_b32
599 ; GCN-DAG: v_accvgpr_read_b32
600 ; GCN-DAG: v_accvgpr_read_b32
601 ; GCN-DAG: v_accvgpr_read_b32
602 ; GCN-DAG: v_accvgpr_read_b32
603 ; GCN-DAG: v_accvgpr_read_b32
604 ; GCN-DAG: v_accvgpr_read_b32
605 ; GCN-DAG: v_accvgpr_read_b32
606 ; GCN-DAG: v_accvgpr_read_b32
607 ; GCN-DAG: v_accvgpr_read_b32
608 ; GCN-DAG: v_accvgpr_read_b32
609 ; GCN-DAG: v_accvgpr_read_b32
610 ; GCN-DAG: v_accvgpr_read_b32
611 ; GCN-DAG: v_accvgpr_read_b32
612 ; GCN-DAG: v_accvgpr_read_b32
613 ; GCN-DAG: v_accvgpr_read_b32
614 ; GCN-DAG: global_store_dwordx4
615 ; GCN-DAG: global_store_dwordx4
616 ; GCN-DAG: global_store_dwordx4
617 ; GCN-DAG: global_store_dwordx4
626 ; GCN-LABEL: {{^}}test_mfma_i32_4x4x4i8:
627 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2
628 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1
629 ; GCN: s_load_dwordx4
630 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
631 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
632 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
633 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
634 ; GCN: v_mfma_i32_4x4x4i8 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:…
635 ; GCN: v_accvgpr_read_b32
636 ; GCN: v_accvgpr_read_b32
637 ; GCN: v_accvgpr_read_b32
638 ; GCN: v_accvgpr_read_b32
639 ; GCN: global_store_dwordx4
648 ; GCN-LABEL: {{^}}test_mfma_i32_32x32x8i8:
649 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2
650 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1
651 ; GCN: s_load_dwordx16
652 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
653 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
654 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
655 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
656 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
657 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
658 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
659 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
660 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
661 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
662 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
663 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
664 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
665 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
666 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
667 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
668 ; GCN: v_mfma_i32_32x32x8i8 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abi…
669 ; GCN-DAG: v_accvgpr_read_b32
670 ; GCN-DAG: v_accvgpr_read_b32
671 ; GCN-DAG: v_accvgpr_read_b32
672 ; GCN-DAG: v_accvgpr_read_b32
673 ; GCN-DAG: v_accvgpr_read_b32
674 ; GCN-DAG: v_accvgpr_read_b32
675 ; GCN-DAG: v_accvgpr_read_b32
676 ; GCN-DAG: v_accvgpr_read_b32
677 ; GCN-DAG: v_accvgpr_read_b32
678 ; GCN-DAG: v_accvgpr_read_b32
679 ; GCN-DAG: v_accvgpr_read_b32
680 ; GCN-DAG: v_accvgpr_read_b32
681 ; GCN-DAG: v_accvgpr_read_b32
682 ; GCN-DAG: v_accvgpr_read_b32
683 ; GCN-DAG: v_accvgpr_read_b32
684 ; GCN-DAG: v_accvgpr_read_b32
685 ; GCN-DAG: global_store_dwordx4
686 ; GCN-DAG: global_store_dwordx4
687 ; GCN-DAG: global_store_dwordx4
688 ; GCN-DAG: global_store_dwordx4
697 ; GCN-LABEL: {{^}}test_mfma_i32_16x16x16i8:
698 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2
699 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1
700 ; GCN: s_load_dwordx4
701 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
702 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
703 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
704 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
705 ; GCN: v_mfma_i32_16x16x16i8 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 ab…
706 ; GCN-DAG: v_accvgpr_read_b32
707 ; GCN-DAG: v_accvgpr_read_b32
708 ; GCN-DAG: v_accvgpr_read_b32
709 ; GCN-DAG: v_accvgpr_read_b32
710 ; GCN-DAG: global_store_dwordx4
719 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x2bf16:
720 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2
721 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1
722 ; GCN-DAG: s_load_dwordx16
723 ; GCN-DAG: s_load_dwordx16
724 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
725 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
726 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
727 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
728 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
729 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
730 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
731 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
732 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
733 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
734 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
735 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
736 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
737 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
738 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
739 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
740 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
741 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
742 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
743 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
744 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
745 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
746 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
747 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
748 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
749 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
750 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
751 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
752 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
753 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
754 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
755 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
756 ; GCN: v_mfma_f32_32x32x2bf16 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 a…
757 ; GCN-DAG: v_accvgpr_read_b32
758 ; GCN-DAG: v_accvgpr_read_b32
759 ; GCN-DAG: v_accvgpr_read_b32
760 ; GCN-DAG: v_accvgpr_read_b32
761 ; GCN-DAG: v_accvgpr_read_b32
762 ; GCN-DAG: v_accvgpr_read_b32
763 ; GCN-DAG: v_accvgpr_read_b32
764 ; GCN-DAG: v_accvgpr_read_b32
765 ; GCN-DAG: v_accvgpr_read_b32
766 ; GCN-DAG: v_accvgpr_read_b32
767 ; GCN-DAG: v_accvgpr_read_b32
768 ; GCN-DAG: v_accvgpr_read_b32
769 ; GCN-DAG: v_accvgpr_read_b32
770 ; GCN-DAG: v_accvgpr_read_b32
771 ; GCN-DAG: v_accvgpr_read_b32
772 ; GCN-DAG: v_accvgpr_read_b32
773 ; GCN-DAG: v_accvgpr_read_b32
774 ; GCN-DAG: v_accvgpr_read_b32
775 ; GCN-DAG: v_accvgpr_read_b32
776 ; GCN-DAG: v_accvgpr_read_b32
777 ; GCN-DAG: v_accvgpr_read_b32
778 ; GCN-DAG: v_accvgpr_read_b32
779 ; GCN-DAG: v_accvgpr_read_b32
780 ; GCN-DAG: v_accvgpr_read_b32
781 ; GCN-DAG: v_accvgpr_read_b32
782 ; GCN-DAG: v_accvgpr_read_b32
783 ; GCN-DAG: v_accvgpr_read_b32
784 ; GCN-DAG: v_accvgpr_read_b32
785 ; GCN-DAG: v_accvgpr_read_b32
786 ; GCN-DAG: v_accvgpr_read_b32
787 ; GCN-DAG: v_accvgpr_read_b32
788 ; GCN-DAG: v_accvgpr_read_b32
789 ; GCN-DAG: global_store_dwordx4
790 ; GCN-DAG: global_store_dwordx4
791 ; GCN-DAG: global_store_dwordx4
792 ; GCN-DAG: global_store_dwordx4
793 ; GCN-DAG: global_store_dwordx4
794 ; GCN-DAG: global_store_dwordx4
795 ; GCN-DAG: global_store_dwordx4
796 ; GCN-DAG: global_store_dwordx4
807 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x2bf16:
808 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2
809 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1
810 ; GCN: s_load_dwordx16
811 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
812 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
813 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
814 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
815 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
816 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
817 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
818 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
819 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
820 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
821 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
822 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
823 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
824 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
825 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
826 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
827 ; GCN: v_mfma_f32_16x16x2bf16 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 a…
828 ; GCN-DAG: v_accvgpr_read_b32
829 ; GCN-DAG: v_accvgpr_read_b32
830 ; GCN-DAG: v_accvgpr_read_b32
831 ; GCN-DAG: v_accvgpr_read_b32
832 ; GCN-DAG: v_accvgpr_read_b32
833 ; GCN-DAG: v_accvgpr_read_b32
834 ; GCN-DAG: v_accvgpr_read_b32
835 ; GCN-DAG: v_accvgpr_read_b32
836 ; GCN-DAG: v_accvgpr_read_b32
837 ; GCN-DAG: v_accvgpr_read_b32
838 ; GCN-DAG: v_accvgpr_read_b32
839 ; GCN-DAG: v_accvgpr_read_b32
840 ; GCN-DAG: v_accvgpr_read_b32
841 ; GCN-DAG: v_accvgpr_read_b32
842 ; GCN-DAG: v_accvgpr_read_b32
843 ; GCN-DAG: v_accvgpr_read_b32
844 ; GCN-DAG: global_store_dwordx4
845 ; GCN-DAG: global_store_dwordx4
846 ; GCN-DAG: global_store_dwordx4
847 ; GCN-DAG: global_store_dwordx4
858 ; GCN-LABEL: {{^}}test_mfma_f32_4x4x2bf16:
859 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2
860 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1
861 ; GCN: s_load_dwordx4
862 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
863 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
864 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
865 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
866 ; GCN: v_mfma_f32_4x4x2bf16 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abi…
867 ; GCN-DAG: v_accvgpr_read_b32
868 ; GCN-DAG: v_accvgpr_read_b32
869 ; GCN-DAG: v_accvgpr_read_b32
870 ; GCN-DAG: v_accvgpr_read_b32
871 ; GCN-DAG: global_store_dwordx4
882 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x4bf16:
883 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2
884 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1
885 ; GCN: s_load_dwordx16
886 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
887 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
888 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
889 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
890 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
891 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
892 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
893 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
894 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
895 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
896 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
897 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
898 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
899 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
900 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
901 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
902 ; GCN: v_mfma_f32_32x32x4bf16 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 a…
903 ; GCN-DAG: v_accvgpr_read_b32
904 ; GCN-DAG: v_accvgpr_read_b32
905 ; GCN-DAG: v_accvgpr_read_b32
906 ; GCN-DAG: v_accvgpr_read_b32
907 ; GCN-DAG: v_accvgpr_read_b32
908 ; GCN-DAG: v_accvgpr_read_b32
909 ; GCN-DAG: v_accvgpr_read_b32
910 ; GCN-DAG: v_accvgpr_read_b32
911 ; GCN-DAG: v_accvgpr_read_b32
912 ; GCN-DAG: v_accvgpr_read_b32
913 ; GCN-DAG: v_accvgpr_read_b32
914 ; GCN-DAG: v_accvgpr_read_b32
915 ; GCN-DAG: v_accvgpr_read_b32
916 ; GCN-DAG: v_accvgpr_read_b32
917 ; GCN-DAG: v_accvgpr_read_b32
918 ; GCN-DAG: v_accvgpr_read_b32
919 ; GCN-DAG: global_store_dwordx4
920 ; GCN-DAG: global_store_dwordx4
921 ; GCN-DAG: global_store_dwordx4
922 ; GCN-DAG: global_store_dwordx4
933 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x8bf16:
934 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2
935 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1
936 ; GCN: s_load_dwordx4
937 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
938 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
939 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
940 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
941 ; GCN: v_mfma_f32_16x16x8bf16 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 a…
942 ; GCN-DAG: v_accvgpr_read_b32
943 ; GCN-DAG: v_accvgpr_read_b32
944 ; GCN-DAG: v_accvgpr_read_b32
945 ; GCN-DAG: v_accvgpr_read_b32
946 ; GCN-DAG: global_store_dwordx4
957 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_forward_acc:
958 ; GCN:      v_mfma_f32_32x32x1f32 [[MAI1:a\[[0-9]+:[0-9]+\]]], v{{[0-9]+}}, v{{[0-9]+}}, a[{{[0-9]+…
959 ; GCN-NEXT: v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v{{[0-9]+}}, [[MAI1]]
969 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x1f32_forward_acc:
970 ; GCN:      v_mfma_f32_16x16x1f32 [[MAI1:a\[[0-9]+:[0-9]+\]]], v{{[0-9]+}}, v{{[0-9]+}}, a[{{[0-9]+…
971 ; GCN-NEXT: v_mfma_f32_16x16x1f32 a[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v{{[0-9]+}}, [[MAI1]]
981 ; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32_forward_acc:
982 ; GCN:      v_mfma_f32_4x4x1f32 [[MAI1:a\[[0-9]+:[0-9]+\]]], v{{[0-9]+}}, v{{[0-9]+}}, a[{{[0-9]+:[…
983 ; GCN-NEXT: v_mfma_f32_4x4x1f32 a[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v{{[0-9]+}}, [[MAI1]]
993 ; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32_imm_splat:
994 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
995 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
996 ; GCN: v_mfma_f32_4x4x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], 1.0
997 ; GCN: v_accvgpr_read_b32
998 ; GCN: v_accvgpr_read_b32
999 ; GCN: v_accvgpr_read_b32
1000 ; GCN: v_accvgpr_read_b32
1001 ; GCN: global_store_dwordx4
1009 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x1f32_imm_splat:
1010 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
1011 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
1012 ; GCN: v_mfma_f32_16x16x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], 1.0
1013 ; GCN-DAG: v_accvgpr_read_b32
1014 ; GCN-DAG: v_accvgpr_read_b32
1015 ; GCN-DAG: v_accvgpr_read_b32
1016 ; GCN-DAG: v_accvgpr_read_b32
1017 ; GCN-DAG: v_accvgpr_read_b32
1018 ; GCN-DAG: v_accvgpr_read_b32
1019 ; GCN-DAG: v_accvgpr_read_b32
1020 ; GCN-DAG: v_accvgpr_read_b32
1021 ; GCN-DAG: v_accvgpr_read_b32
1022 ; GCN-DAG: v_accvgpr_read_b32
1023 ; GCN-DAG: v_accvgpr_read_b32
1024 ; GCN-DAG: v_accvgpr_read_b32
1025 ; GCN-DAG: v_accvgpr_read_b32
1026 ; GCN-DAG: v_accvgpr_read_b32
1027 ; GCN-DAG: v_accvgpr_read_b32
1028 ; GCN-DAG: v_accvgpr_read_b32
1029 ; GCN-DAG: global_store_dwordx4
1030 ; GCN-DAG: global_store_dwordx4
1031 ; GCN-DAG: global_store_dwordx4
1032 ; GCN-DAG: global_store_dwordx4
1040 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x8f16_imm_splat:
1041 ; GCN-DAG: v_mov_b32_e32 v[[TWO:[0-9]+]], 0x40004000
1042 ; GCN-DAG: v_mov_b32_e32 v[[ONE:[0-9]+]], 0x3c003c00
1043 ; GCN: v_mfma_f32_32x32x8f16 a[{{[0-9]+:[0-9]+}}], v{{\[}}[[ONE]]:{{[0-9]+}}], v{{\[}}[[TWO]]:{{[0-…
1044 ; GCN-DAG: v_accvgpr_read_b32
1045 ; GCN-DAG: v_accvgpr_read_b32
1046 ; GCN-DAG: v_accvgpr_read_b32
1047 ; GCN-DAG: v_accvgpr_read_b32
1048 ; GCN-DAG: v_accvgpr_read_b32
1049 ; GCN-DAG: v_accvgpr_read_b32
1050 ; GCN-DAG: v_accvgpr_read_b32
1051 ; GCN-DAG: v_accvgpr_read_b32
1052 ; GCN-DAG: v_accvgpr_read_b32
1053 ; GCN-DAG: v_accvgpr_read_b32
1054 ; GCN-DAG: v_accvgpr_read_b32
1055 ; GCN-DAG: v_accvgpr_read_b32
1056 ; GCN-DAG: v_accvgpr_read_b32
1057 ; GCN-DAG: v_accvgpr_read_b32
1058 ; GCN-DAG: v_accvgpr_read_b32
1059 ; GCN-DAG: v_accvgpr_read_b32
1060 ; GCN-DAG: global_store_dwordx4
1061 ; GCN-DAG: global_store_dwordx4
1062 ; GCN-DAG: global_store_dwordx4
1063 ; GCN-DAG: global_store_dwordx4
1071 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_imm_splat:
1072 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
1073 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
1074 ; GCN: v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], 0
1075 ; GCN-DAG: v_accvgpr_read_b32
1076 ; GCN-DAG: v_accvgpr_read_b32
1077 ; GCN-DAG: v_accvgpr_read_b32
1078 ; GCN-DAG: v_accvgpr_read_b32
1079 ; GCN-DAG: v_accvgpr_read_b32
1080 ; GCN-DAG: v_accvgpr_read_b32
1081 ; GCN-DAG: v_accvgpr_read_b32
1082 ; GCN-DAG: v_accvgpr_read_b32
1083 ; GCN-DAG: v_accvgpr_read_b32
1084 ; GCN-DAG: v_accvgpr_read_b32
1085 ; GCN-DAG: v_accvgpr_read_b32
1086 ; GCN-DAG: v_accvgpr_read_b32
1087 ; GCN-DAG: v_accvgpr_read_b32
1088 ; GCN-DAG: v_accvgpr_read_b32
1089 ; GCN-DAG: v_accvgpr_read_b32
1090 ; GCN-DAG: v_accvgpr_read_b32
1091 ; GCN-DAG: v_accvgpr_read_b32
1092 ; GCN-DAG: v_accvgpr_read_b32
1093 ; GCN-DAG: v_accvgpr_read_b32
1094 ; GCN-DAG: v_accvgpr_read_b32
1095 ; GCN-DAG: v_accvgpr_read_b32
1096 ; GCN-DAG: v_accvgpr_read_b32
1097 ; GCN-DAG: v_accvgpr_read_b32
1098 ; GCN-DAG: v_accvgpr_read_b32
1099 ; GCN-DAG: v_accvgpr_read_b32
1100 ; GCN-DAG: v_accvgpr_read_b32
1101 ; GCN-DAG: v_accvgpr_read_b32
1102 ; GCN-DAG: v_accvgpr_read_b32
1103 ; GCN-DAG: v_accvgpr_read_b32
1104 ; GCN-DAG: v_accvgpr_read_b32
1105 ; GCN-DAG: v_accvgpr_read_b32
1106 ; GCN-DAG: v_accvgpr_read_b32
1107 ; GCN-DAG: global_store_dwordx4
1108 ; GCN-DAG: global_store_dwordx4
1109 ; GCN-DAG: global_store_dwordx4
1110 ; GCN-DAG: global_store_dwordx4
1111 ; GCN-DAG: global_store_dwordx4
1112 ; GCN-DAG: global_store_dwordx4
1113 ; GCN-DAG: global_store_dwordx4
1114 ; GCN-DAG: global_store_dwordx4
1122 ; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32_imm:
1123 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1124 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1125 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1126 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 2.0
1127 ; GCN: v_mfma_f32_4x4x1f32 a[{{[0-9]+:[0-9]+}}], {{v[0-9]+}}, {{v[0-9]+}}, a[{{[0-9]+:[0-9]+}}]
1128 ; GCN: v_accvgpr_read_b32
1129 ; GCN: v_accvgpr_read_b32
1130 ; GCN: v_accvgpr_read_b32
1131 ; GCN: v_accvgpr_read_b32
1132 ; GCN: global_store_dwordx4
1140 ; GCN-LABEL: {{^}}test_mfma_f32_16x16x1f32_imm:
1141 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1142 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 2.0
1143 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1144 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1145 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1146 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1147 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1148 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1149 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1150 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1151 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1152 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1153 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1154 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1155 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1156 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1157 ; GCN: v_mfma_f32_16x16x1f32 a[{{[0-9]+:[0-9]+}}], {{v[0-9]+}}, {{v[0-9]+}}, a[{{[0-9]+:[0-9]+}}]
1158 ; GCN-DAG: v_accvgpr_read_b32
1159 ; GCN-DAG: v_accvgpr_read_b32
1160 ; GCN-DAG: v_accvgpr_read_b32
1161 ; GCN-DAG: v_accvgpr_read_b32
1162 ; GCN-DAG: v_accvgpr_read_b32
1163 ; GCN-DAG: v_accvgpr_read_b32
1164 ; GCN-DAG: v_accvgpr_read_b32
1165 ; GCN-DAG: v_accvgpr_read_b32
1166 ; GCN-DAG: v_accvgpr_read_b32
1167 ; GCN-DAG: v_accvgpr_read_b32
1168 ; GCN-DAG: v_accvgpr_read_b32
1169 ; GCN-DAG: v_accvgpr_read_b32
1170 ; GCN-DAG: v_accvgpr_read_b32
1171 ; GCN-DAG: v_accvgpr_read_b32
1172 ; GCN-DAG: v_accvgpr_read_b32
1173 ; GCN-DAG: v_accvgpr_read_b32
1174 ; GCN-DAG: global_store_dwordx4
1175 ; GCN-DAG: global_store_dwordx4
1176 ; GCN-DAG: global_store_dwordx4
1177 ; GCN-DAG: global_store_dwordx4
1185 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_imm:
1186 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1187 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
1188 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1189 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1190 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1191 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1192 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1193 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1194 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1195 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1196 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1197 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1198 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1199 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1200 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1201 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1202 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1203 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1204 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1205 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1206 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1207 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1208 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1209 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1210 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1211 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1212 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1213 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1214 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1215 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1216 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1217 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0
1218 ; GCN: v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], {{v[0-9]+}}, {{v[0-9]+}}, a[{{[0-9]+:[0-9]+}}]
1219 ; GCN-DAG: v_accvgpr_read_b32
1220 ; GCN-DAG: v_accvgpr_read_b32
1221 ; GCN-DAG: v_accvgpr_read_b32
1222 ; GCN-DAG: v_accvgpr_read_b32
1223 ; GCN-DAG: v_accvgpr_read_b32
1224 ; GCN-DAG: v_accvgpr_read_b32
1225 ; GCN-DAG: v_accvgpr_read_b32
1226 ; GCN-DAG: v_accvgpr_read_b32
1227 ; GCN-DAG: v_accvgpr_read_b32
1228 ; GCN-DAG: v_accvgpr_read_b32
1229 ; GCN-DAG: v_accvgpr_read_b32
1230 ; GCN-DAG: v_accvgpr_read_b32
1231 ; GCN-DAG: v_accvgpr_read_b32
1232 ; GCN-DAG: v_accvgpr_read_b32
1233 ; GCN-DAG: v_accvgpr_read_b32
1234 ; GCN-DAG: v_accvgpr_read_b32
1235 ; GCN-DAG: v_accvgpr_read_b32
1236 ; GCN-DAG: v_accvgpr_read_b32
1237 ; GCN-DAG: v_accvgpr_read_b32
1238 ; GCN-DAG: v_accvgpr_read_b32
1239 ; GCN-DAG: v_accvgpr_read_b32
1240 ; GCN-DAG: v_accvgpr_read_b32
1241 ; GCN-DAG: v_accvgpr_read_b32
1242 ; GCN-DAG: v_accvgpr_read_b32
1243 ; GCN-DAG: v_accvgpr_read_b32
1244 ; GCN-DAG: v_accvgpr_read_b32
1245 ; GCN-DAG: v_accvgpr_read_b32
1246 ; GCN-DAG: v_accvgpr_read_b32
1247 ; GCN-DAG: v_accvgpr_read_b32
1248 ; GCN-DAG: v_accvgpr_read_b32
1249 ; GCN-DAG: v_accvgpr_read_b32
1250 ; GCN-DAG: v_accvgpr_read_b32
1251 ; GCN-DAG: global_store_dwordx4
1252 ; GCN-DAG: global_store_dwordx4
1253 ; GCN-DAG: global_store_dwordx4
1254 ; GCN-DAG: global_store_dwordx4
1255 ; GCN-DAG: global_store_dwordx4
1256 ; GCN-DAG: global_store_dwordx4
1257 ; GCN-DAG: global_store_dwordx4
1258 ; GCN-DAG: global_store_dwordx4
1266 ; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32_lit_splat:
1267 ; GCN: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x42f60000
1268 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]
1269 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]
1270 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]
1271 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]
1272 ; GCN: v_mfma_f32_4x4x1f32 a[{{[0-9]+:[0-9]+}}], {{v[0-9]+}}, {{v[0-9]+}}, a[{{[0-9]+:[0-9]+}}]
1273 ; GCN: v_accvgpr_read_b32
1274 ; GCN: v_accvgpr_read_b32
1275 ; GCN: v_accvgpr_read_b32
1276 ; GCN: v_accvgpr_read_b32
1277 ; GCN: global_store_dwordx4
1285 ; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_vecarg:
1286 ; GCN-DAG: v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
1287 ; GCN-DAG: v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
1296 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1297 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1298 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1299 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1300 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1301 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1302 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1303 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1304 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1305 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1306 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1307 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1308 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1309 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1310 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1311 ; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
1312 ; GCN: v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 ab…
1313 ; GCN-DAG: v_accvgpr_read_b32
1314 ; GCN-DAG: v_accvgpr_read_b32
1315 ; GCN-DAG: v_accvgpr_read_b32
1316 ; GCN-DAG: v_accvgpr_read_b32
1317 ; GCN-DAG: v_accvgpr_read_b32
1318 ; GCN-DAG: v_accvgpr_read_b32
1319 ; GCN-DAG: v_accvgpr_read_b32
1320 ; GCN-DAG: v_accvgpr_read_b32
1321 ; GCN-DAG: v_accvgpr_read_b32
1322 ; GCN-DAG: v_accvgpr_read_b32
1323 ; GCN-DAG: v_accvgpr_read_b32
1324 ; GCN-DAG: v_accvgpr_read_b32
1325 ; GCN-DAG: v_accvgpr_read_b32
1326 ; GCN-DAG: v_accvgpr_read_b32
1327 ; GCN-DAG: v_accvgpr_read_b32
1328 ; GCN-DAG: v_accvgpr_read_b32
1329 ; GCN-DAG: v_accvgpr_read_b32
1330 ; GCN-DAG: v_accvgpr_read_b32
1331 ; GCN-DAG: v_accvgpr_read_b32
1332 ; GCN-DAG: v_accvgpr_read_b32
1333 ; GCN-DAG: v_accvgpr_read_b32
1334 ; GCN-DAG: v_accvgpr_read_b32
1335 ; GCN-DAG: v_accvgpr_read_b32
1336 ; GCN-DAG: v_accvgpr_read_b32
1337 ; GCN-DAG: v_accvgpr_read_b32
1338 ; GCN-DAG: v_accvgpr_read_b32
1339 ; GCN-DAG: v_accvgpr_read_b32
1340 ; GCN-DAG: v_accvgpr_read_b32
1341 ; GCN-DAG: v_accvgpr_read_b32
1342 ; GCN-DAG: v_accvgpr_read_b32
1343 ; GCN-DAG: v_accvgpr_read_b32
1344 ; GCN-DAG: v_accvgpr_read_b32
1345 ; GCN-DAG: global_store_dwordx4
1346 ; GCN-DAG: global_store_dwordx4
1347 ; GCN-DAG: global_store_dwordx4
1348 ; GCN-DAG: global_store_dwordx4
1349 ; GCN-DAG: global_store_dwordx4
1350 ; GCN-DAG: global_store_dwordx4
1351 ; GCN-DAG: global_store_dwordx4
1352 ; GCN-DAG: global_store_dwordx4