1; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN,NOLIT-SRCC,GFX908,GFX908_A %s
2; RUN: llc -march=amdgcn -mcpu=gfx908 -mattr=-mfma-inline-literal-bug -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN,LIT-SRCC,GFX908,GFX908_A %s
3; RUN: llc -march=amdgcn -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX90A,GFX908_A %s
4
5declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x float>, i32, i32, i32)
6declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float, float, <16 x float>, i32, i32, i32)
7declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32, i32, i32)
8declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x2f32(float, float, <16 x float>, i32, i32, i32)
9declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x4f32(float, float, <4 x float>, i32, i32, i32)
10declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x4f16(<4 x half>, <4 x half>, <32 x float>, i32, i32, i32)
11declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x4f16(<4 x half>, <4 x half>, <16 x float>, i32, i32, i32)
12declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x4f16(<4 x half>, <4 x half>, <4 x float>, i32, i32, i32)
13declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half>, <4 x half>, <16 x float>, i32, i32, i32)
14declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half>, <4 x half>, <4 x float>, i32, i32, i32)
15declare <32 x i32> @llvm.amdgcn.mfma.i32.32x32x4i8(i32, i32, <32 x i32>, i32, i32, i32)
16declare <16 x i32> @llvm.amdgcn.mfma.i32.16x16x4i8(i32, i32, <16 x i32>, i32, i32, i32)
17declare <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32, i32, <4 x i32>, i32, i32, i32)
18declare i32 @llvm.amdgcn.workitem.id.x()
19
20; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32:
21; GCN-DAG:        v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
22; GCN-DAG:        v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
23; GCN-DAG:        s_load_dwordx16
24; GCN-DAG:        s_load_dwordx16
25; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
26; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
27; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
28; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
29; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
30; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
31; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
32; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
33; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
34; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
35; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
36; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
37; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
38; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
39; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
40; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
41; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
42; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
43; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
44; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
45; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
46; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
47; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
48; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
49; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
50; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
51; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
52; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
53; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
54; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
55; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
56; GFX908_A-DAG:   v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
57; GFX908_A:       v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
58; GFX908-COUNT-4: v_accvgpr_read_b32
59; GFX908-COUNT-2: global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}]
60; GFX908-COUNT-4: v_accvgpr_read_b32
61; GFX908-COUNT-2: global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}]
62; GFX908-COUNT-4: v_accvgpr_read_b32
63; GFX908-COUNT-2: global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}]
64; GFX908-COUNT-4: v_accvgpr_read_b32
65; GFX908-COUNT-2: global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}]
66; GFX90A-NOT:     v_accvgpr_read_b32
67; GFX90A-COUNT-8: global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}]
68define amdgpu_kernel void @test_mfma_f32_32x32x1f32(<32 x float> addrspace(1)* %arg) {
69bb:
70  %in.1 = load <32 x float>, <32 x float> addrspace(1)* %arg
71  %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 1, i32 2, i32 3)
72  store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
73  ret void
74}
75
76; GCN-LABEL: {{^}}test_mfma_f32_16x16x1f32:
77; GCN-DAG:           v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
78; GCN-DAG:           v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
79; GCN:               s_load_dwordx16
80; GFX908_A-COUNT-16: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
81; GFX908_A:          v_mfma_f32_16x16x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
82; GFX908-COUNT:      v_accvgpr_read_b32
83; GFX908-COUNT-4:    global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}]
84; GFX90A-NOT:        v_accvgpr_read_b32
85; GFX90A-COUNT-4:    global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}]
86define amdgpu_kernel void @test_mfma_f32_16x16x1f32(<16 x float> addrspace(1)* %arg) {
87bb:
88  %in.1 = load <16 x float>, <16 x float> addrspace(1)* %arg
89  %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float 1.0, float 2.0, <16 x float> %in.1, i32 1, i32 2, i32 3)
90  store <16 x float> %mai.1, <16 x float> addrspace(1)* %arg
91  ret void
92}
93
94; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32:
95; GCN-DAG:          v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
96; GCN-DAG:          v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
97; GCN:              s_load_dwordx4
98; GFX908_A-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
99; GFX908_A:         v_mfma_f32_4x4x1f32 [[RES:a\[[0-9]+:[0-9]+\]]], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
100; GFX908-COUNT-4:   v_accvgpr_read_b32
101; GFX908:           global_store_dwordx4
102; GFX90A-NOT:       v_accvgpr_read_b32
103; GFX90A:           global_store_dwordx4 {{v[0-9]+}}, [[RES]]
104define amdgpu_kernel void @test_mfma_f32_4x4x1f32(<4 x float> addrspace(1)* %arg) {
105bb:
106  %in.1 = load <4 x float>, <4 x float> addrspace(1)* %arg
107  %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %in.1, i32 1, i32 2, i32 3)
108  store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
109  ret void
110}
111
112; GCN-LABEL: {{^}}test_mfma_f32_32x32x2f32:
113; GCN-DAG:           v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
114; GCN-DAG:           v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
115; GCN:               s_load_dwordx16
116; GFX908_A-COUNT-16: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
117; GFX908_A:          v_mfma_f32_32x32x2f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
118; GFX908-COUNT-16:   v_accvgpr_read_b32
119; GFX908-COUNT-4:    global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}]
120; GFX90A-NOT:        v_accvgpr_read_b32
121; GFX90A-COUNT-4:    global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}]
122define amdgpu_kernel void @test_mfma_f32_32x32x2f32(<16 x float> addrspace(1)* %arg) {
123bb:
124  %in.1 = load <16 x float>, <16 x float> addrspace(1)* %arg
125  %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x2f32(float 1.0, float 2.0, <16 x float> %in.1, i32 1, i32 2, i32 3)
126  store <16 x float> %mai.1, <16 x float> addrspace(1)* %arg
127  ret void
128}
129
130; GCN-LABEL: {{^}}test_mfma_f32_16x16x4f32:
131; GCN-DAG:          v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
132; GCN-DAG:          v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
133; GCN:              s_load_dwordx4
134; GFX908_A-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
135; GFX908_A:         v_mfma_f32_16x16x4f32 [[RES:a\[[0-9]+:[0-9]+\]]], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
136; GFX908-COUNT-4:   v_accvgpr_read_b32
137; GFX908:           global_store_dwordx4
138; GFX90A-NOT:       v_accvgpr_read_b32
139; GFX90A:           global_store_dwordx4 {{v[0-9]+}}, [[RES]],
140define amdgpu_kernel void @test_mfma_f32_16x16x4f32(<4 x float> addrspace(1)* %arg) {
141bb:
142  %in.1 = load <4 x float>, <4 x float> addrspace(1)* %arg
143  %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x4f32(float 1.0, float 2.0, <4 x float> %in.1, i32 1, i32 2, i32 3)
144  store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
145  ret void
146}
147
148; GCN-LABEL: {{^}}test_mfma_f32_32x32x4f16:
149; GCN-DAG:           s_load_dwordx16
150; GCN-DAG:           s_load_dwordx16
151; GFX908_A-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
152; GFX908_A:          v_mfma_f32_32x32x4f16 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
153; GFX908-COUNT-32:   v_accvgpr_read_b32
154; GFX908:            global_store_dwordx4
155; GFX90A-NOT:        v_accvgpr_read_b32
156; GFX90A-COUNT-8:    global_store_dwordx4 {{v[0-9]+}}, a[{{[0-9:]+}}],
157define amdgpu_kernel void @test_mfma_f32_32x32x4f16(<32 x float> addrspace(1)* %arg, <4 x half> addrspace(1)* %c) {
158bb:
159  %in.1 = load <32 x float>, <32 x float> addrspace(1)* %arg
160  %c.1 = load <4 x half>, <4 x half> addrspace(1)* %c
161  %c2p = getelementptr <4 x half>, <4 x half> addrspace(1)* %c, i64 1
162  %c.2 = load <4 x half>, <4 x half> addrspace(1)* %c2p
163  %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x4f16(<4 x half> %c.1, <4 x half> %c.2, <32 x float> %in.1, i32 1, i32 2, i32 3)
164  store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
165  ret void
166}
167
168; GCN-LABEL: {{^}}test_mfma_f32_16x16x4f16:
169; GCN:               s_load_dwordx16
170; GFX908_A-COUNT-16: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
171; GFX908_A:          v_mfma_f32_16x16x4f16 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
172; GFX908-COUNT-16:   v_accvgpr_read_b32
173; GFX908:            global_store_dwordx4
174; GFX90A-NOT:        v_accvgpr_read_b32
175; GFX90A-COUNT-4:    global_store_dwordx4 {{v[0-9]+}}, a[{{[0-9:]+}}],
176define amdgpu_kernel void @test_mfma_f32_16x16x4f16(<16 x float> addrspace(1)* %arg, <4 x half> addrspace(1)* %c) {
177bb:
178  %in.1 = load <16 x float>, <16 x float> addrspace(1)* %arg
179  %c.1 = load <4 x half>, <4 x half> addrspace(1)* %c
180  %c2p = getelementptr <4 x half>, <4 x half> addrspace(1)* %c, i64 1
181  %c.2 = load <4 x half>, <4 x half> addrspace(1)* %c2p
182  %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x4f16(<4 x half> %c.1, <4 x half> %c.2, <16 x float> %in.1, i32 1, i32 2, i32 3)
183  store <16 x float> %mai.1, <16 x float> addrspace(1)* %arg
184  ret void
185}
186
187; GCN-LABEL: {{^}}test_mfma_f32_4x4x4f16:
188; GCN:              s_load_dwordx4
189; GCN:              s_load_dwordx2
190; GCN:              s_load_dwordx2
191; GFX908_A-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
192; GFX908_A:         v_mfma_f32_4x4x4f16 [[RES: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
193; GFX908-COUNT-4:   v_accvgpr_read_b32
194; GFX908:           global_store_dwordx4
195; GFX90A-NOT:       v_accvgpr_read_b32
196; GFX90A:           global_store_dwordx4 {{v[0-9]+}}, [[RES]],
197define amdgpu_kernel void @test_mfma_f32_4x4x4f16(<4 x float> addrspace(1)* %arg, <4 x half> addrspace(1)* %c) {
198bb:
199  %in.1 = load <4 x float>, <4 x float> addrspace(1)* %arg
200  %c.1 = load <4 x half>, <4 x half> addrspace(1)* %c
201  %c2p = getelementptr <4 x half>, <4 x half> addrspace(1)* %c, i64 1
202  %c.2 = load <4 x half>, <4 x half> addrspace(1)* %c2p
203  %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x4f16(<4 x half> %c.1, <4 x half> %c.2, <4 x float> %in.1, i32 1, i32 2, i32 3)
204  store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
205  ret void
206}
207
208; GCN-LABEL: {{^}}test_mfma_f32_32x32x8f16:
209; GCN:               s_load_dwordx16
210; GCN:               s_waitcnt lgkmcnt(0)
211; GFX908_A:          v_mov_b32_e32 v{{[0-9]+}}, s{{[0-9]+}}
212; GFX908_A-COUNT-16: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
213; GFX908_A:          v_mfma_f32_32x32x8f16 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
214; GFX908-COUNT-16:   v_accvgpr_read_b32
215; GFX908:            global_store_dwordx4
216; GFX90A-NOT:        v_accvgpr_read_b32
217; GFX90A-COUNT-4:    global_store_dwordx4 {{v[0-9]+}}, a[{{[0-9:]+}}],
218define amdgpu_kernel void @test_mfma_f32_32x32x8f16(<16 x float> addrspace(1)* %arg, <4 x half> addrspace(1)* %c) {
219bb:
220  %in.1 = load <16 x float>, <16 x float> addrspace(1)* %arg
221  %c.1 = load <4 x half>, <4 x half> addrspace(1)* %c
222  %c2p = getelementptr <4 x half>, <4 x half> addrspace(1)* %c, i64 1
223  %c.2 = load <4 x half>, <4 x half> addrspace(1)* %c2p
224  %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> %c.1, <4 x half> %c.2, <16 x float> %in.1, i32 1, i32 2, i32 3)
225  store <16 x float> %mai.1, <16 x float> addrspace(1)* %arg
226  ret void
227}
228
229; GCN-LABEL: {{^}}test_mfma_f32_16x16x16f16:
230; GCN:              s_load_dwordx4
231; GCN:              s_load_dwordx4
232; GFX908_A-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
233; GFX908_A:         v_mfma_f32_16x16x16f16 [[RES: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
234; GFX908-COUNT-4:   v_accvgpr_read_b32
235; GFX908:           global_store_dwordx4
236; GFX90A-NOT:       v_accvgpr_read_b32
237; GFX90A:           global_store_dwordx4 {{v[0-9]+}}, [[RES]],
238define amdgpu_kernel void @test_mfma_f32_16x16x16f16(<4 x float> addrspace(1)* %arg, <4 x half> addrspace(1)* %c) {
239bb:
240  %in.1 = load <4 x float>, <4 x float> addrspace(1)* %arg
241  %c.1 = load <4 x half>, <4 x half> addrspace(1)* %c
242  %c2p = getelementptr <4 x half>, <4 x half> addrspace(1)* %c, i64 1
243  %c.2 = load <4 x half>, <4 x half> addrspace(1)* %c2p
244  %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> %c.1, <4 x half> %c.2, <4 x float> %in.1, i32 1, i32 2, i32 3)
245  store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
246  ret void
247}
248
249; GCN-LABEL: {{^}}test_mfma_i32_32x32x4i8:
250; GCN-DAG:         v_mov_b32_e32 [[TWO:v[0-9]+]], 2
251; GCN-DAG:         v_mov_b32_e32 [[ONE:v[0-9]+]], 1
252; GCN-DAG:         s_load_dwordx16
253; GCN-DAG:         s_load_dwordx16
254; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
255; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
256; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
257; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
258; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
259; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
260; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
261; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
262; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
263; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
264; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
265; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
266; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
267; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
268; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
269; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
270; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
271; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
272; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
273; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
274; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
275; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
276; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
277; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
278; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
279; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
280; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
281; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
282; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
283; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
284; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
285; GFX908_A-DAG:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
286; GFX908_A:        v_mfma_i32_32x32x4i8 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
287; GFX908-COUNT-32: v_accvgpr_read_b32
288; GFX908:          global_store_dwordx4
289; GFX90A-NOT:      v_accvgpr_read_b32
290; GFX90A-COUNT-8:  global_store_dwordx4 {{v[0-9]+}}, a[{{[0-9:]+}}],
291define amdgpu_kernel void @test_mfma_i32_32x32x4i8(<32 x i32> addrspace(1)* %arg) {
292bb:
293  %in.1 = load <32 x i32>, <32 x i32> addrspace(1)* %arg
294  %mai.1 = tail call <32 x i32> @llvm.amdgcn.mfma.i32.32x32x4i8(i32 1, i32 2, <32 x i32> %in.1, i32 1, i32 2, i32 3)
295  store <32 x i32> %mai.1, <32 x i32> addrspace(1)* %arg
296  ret void
297}
298
299; GCN-LABEL: {{^}}test_mfma_i32_16x16x4i8:
300; GCN-DAG:           v_mov_b32_e32 [[TWO:v[0-9]+]], 2
301; GCN-DAG:           v_mov_b32_e32 [[ONE:v[0-9]+]], 1
302; GCN:               s_load_dwordx16
303; GFX908_A-COUNT-16: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
304; GFX908_A:          v_mfma_i32_16x16x4i8 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
305; GFX908-COUNT-16:   v_accvgpr_read_b32
306; GFX908:            global_store_dwordx4
307; GFX90A-NOT:        v_accvgpr_read_b32
308; GFX90A-COUNT-4:    global_store_dwordx4 {{v[0-9]+}}, a[{{[0-9:]+}}],
309define amdgpu_kernel void @test_mfma_i32_16x16x4i8(<16 x i32> addrspace(1)* %arg) {
310bb:
311  %in.1 = load <16 x i32>, <16 x i32> addrspace(1)* %arg
312  %mai.1 = tail call <16 x i32> @llvm.amdgcn.mfma.i32.16x16x4i8(i32 1, i32 2, <16 x i32> %in.1, i32 1, i32 2, i32 3)
313  store <16 x i32> %mai.1, <16 x i32> addrspace(1)* %arg
314  ret void
315}
316
317; GCN-LABEL: {{^}}test_mfma_i32_4x4x4i8:
318; GCN-DAG:          v_mov_b32_e32 [[TWO:v[0-9]+]], 2
319; GCN-DAG:          v_mov_b32_e32 [[ONE:v[0-9]+]], 1
320; GCN:              s_load_dwordx4
321; GFX908_A-COUNT-4: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
322; GFX908_A:         v_mfma_i32_4x4x4i8 [[RES:a\[[0-9]+:[0-9]+\]]], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
323; GFX908-COUNT-4:   v_accvgpr_read_b32
324; GFX908:           global_store_dwordx4
325; GFX90A-NOT:       v_accvgpr_read_b32
326; GFX90A:           global_store_dwordx4 {{v[0-9]+}}, [[RES]],
327define amdgpu_kernel void @test_mfma_i32_4x4x4i8(<4 x i32> addrspace(1)* %arg) {
328bb:
329  %in.1 = load <4 x i32>, <4 x i32> addrspace(1)* %arg
330  %mai.1 = tail call <4 x i32> @llvm.amdgcn.mfma.i32.4x4x4i8(i32 1, i32 2, <4 x i32> %in.1, i32 1, i32 2, i32 3)
331  store <4 x i32> %mai.1, <4 x i32> addrspace(1)* %arg
332  ret void
333}
334
335; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_forward_acc:
336; GFX908_A:      v_mfma_f32_32x32x1f32 [[MAI1:a\[[0-9]+:[0-9]+\]]], v{{[0-9]+}}, v{{[0-9]+}}, a[{{[0-9]+:[0-9]+}}]
337; GFX908_A-NEXT: v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v{{[0-9]+}}, [[MAI1]]
338define amdgpu_kernel void @test_mfma_f32_32x32x1f32_forward_acc(<32 x float> addrspace(1)* %arg) {
339bb:
340  %in.1 = load <32 x float>, <32 x float> addrspace(1)* %arg
341  %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 0, i32 0, i32 0)
342  %mai.2 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %mai.1, i32 0, i32 0, i32 0)
343  store <32 x float> %mai.2, <32 x float> addrspace(1)* %arg
344  ret void
345}
346
347; GCN-LABEL: {{^}}test_mfma_f32_16x16x1f32_forward_acc:
348; GFX908_A:      v_mfma_f32_16x16x1f32 [[MAI1:a\[[0-9]+:[0-9]+\]]], v{{[0-9]+}}, v{{[0-9]+}}, a[{{[0-9]+:[0-9]+}}]
349; GFX908_A-NEXT: v_mfma_f32_16x16x1f32 a[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v{{[0-9]+}}, [[MAI1]]
350define amdgpu_kernel void @test_mfma_f32_16x16x1f32_forward_acc(<16 x float> addrspace(1)* %arg) {
351bb:
352  %in.1 = load <16 x float>, <16 x float> addrspace(1)* %arg
353  %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float 1.0, float 2.0, <16 x float> %in.1, i32 0, i32 0, i32 0)
354  %mai.2 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float 1.0, float 2.0, <16 x float> %mai.1, i32 0, i32 0, i32 0)
355  store <16 x float> %mai.2, <16 x float> addrspace(1)* %arg
356  ret void
357}
358
359; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32_forward_acc:
360; GFX908_A:      v_mfma_f32_4x4x1f32 [[MAI1:a\[[0-9]+:[0-9]+\]]], v{{[0-9]+}}, v{{[0-9]+}}, a[{{[0-9]+:[0-9]+}}]
361; GFX908_A-NEXT: v_mfma_f32_4x4x1f32 a[{{[0-9]+:[0-9]+}}], v{{[0-9]+}}, v{{[0-9]+}}, [[MAI1]]
362define amdgpu_kernel void @test_mfma_f32_4x4x1f32_forward_acc(<4 x float> addrspace(1)* %arg) {
363bb:
364  %in.1 = load <4 x float>, <4 x float> addrspace(1)* %arg
365  %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %in.1, i32 0, i32 0, i32 0)
366  %mai.2 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %mai.1, i32 0, i32 0, i32 0)
367  store <4 x float> %mai.2, <4 x float> addrspace(1)* %arg
368  ret void
369}
370
371; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32_imm_splat:
372; GCN-DAG:        v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
373; GCN-DAG:        v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
374; NOLIT-SRCC-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
375; NOLIT-SRCC-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
376; NOLIT-SRCC-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
377; NOLIT-SRCC-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
378; NOLIT-SRCC:     v_mfma_f32_4x4x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9:]+}}]
379; LIT-SRCC:       v_mfma_f32_4x4x1f32 [[RES:a\[[0-9]+:[0-9]+\]]], [[ONE]], [[TWO]], 1.0
380; GFX90A:         v_mfma_f32_4x4x1f32 [[RES:a\[[0-9]+:[0-9]+\]]], [[ONE]], [[TWO]], 1.0
381; GFX908-COUNT-4: v_accvgpr_read_b32
382; GFX908:         global_store_dwordx4
383; GFX90A-NOT:     v_accvgpr_read_b32
384; GFX90A:         global_store_dwordx4 {{v[0-9]+}}, [[RES]],
385define amdgpu_kernel void @test_mfma_f32_4x4x1f32_imm_splat(<4 x float> addrspace(1)* %arg) {
386bb:
387  %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> <float 1.0, float 1.0, float 1.0, float 1.0>, i32 0, i32 0, i32 0)
388  store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
389  ret void
390}
391
392; GCN-LABEL: {{^}}test_mfma_f32_16x16x1f32_imm_splat:
393; GCN-DAG:         v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
394; GCN-DAG:         v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
395; NOLIT-SRCC-DAG:  v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
396; NOLIT-SRCC:      v_mfma_f32_16x16x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9:]+}}]
397; LIT-SRCC:        v_mfma_f32_16x16x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], 1.0
398; GFX90A:          v_mfma_f32_16x16x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], 1.0
399; GFX908-COUNT-16: v_accvgpr_read_b32
400; GFX908:          global_store_dwordx4
401; GFX90A-NOT:      v_accvgpr_read_b32
402; GFX90A-COUNT-4:  global_store_dwordx4 {{v[0-9]+}}, a[{{[0-9:]+}}],
403define amdgpu_kernel void @test_mfma_f32_16x16x1f32_imm_splat(<16 x float> addrspace(1)* %arg) {
404bb:
405  %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float 1.0, float 2.0, <16 x float> <float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0>, i32 0, i32 0, i32 0)
406  store <16 x float> %mai.1, <16 x float> addrspace(1)* %arg
407  ret void
408}
409
410; GCN-LABEL: {{^}}test_mfma_f32_32x32x8f16_imm_splat:
411; GCN-DAG:         v_mov_b32_e32 v[[TWO:[0-9]+]], 0x40004000
412; GCN-DAG:         v_mov_b32_e32 v[[ONE:[0-9]+]], 0x3c003c00
413; NOLIT-SRCC-DAG:  v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
414; NOLIT-SRCC:      v_mfma_f32_32x32x8f16 a[{{[0-9]+:[0-9]+}}], v{{\[}}[[ONE]]:{{[0-9]+}}], v{{\[}}[[TWO]]:{{[0-9]+}}], a[{{[0-9:]+}}]
415; LIT-SRCC:        v_mfma_f32_32x32x8f16 a[{{[0-9]+:[0-9]+}}], v{{\[}}[[ONE]]:{{[0-9]+}}], v{{\[}}[[TWO]]:{{[0-9]+}}], 1.0
416; GFX90A:          v_mfma_f32_32x32x8f16 a[{{[0-9]+:[0-9]+}}], v{{\[}}[[ONE]]:{{[0-9]+}}], v{{\[}}[[TWO]]:{{[0-9]+}}], 1.0
417; GFX908-COUNT-16: v_accvgpr_read_b32
418; GFX908:          global_store_dwordx4
419; GFX90A-NOT:      v_accvgpr_read_b32
420; GFX90A-COUNT-4:  global_store_dwordx4 {{v[0-9]+}}, a[{{[0-9:]+}}],
421define amdgpu_kernel void @test_mfma_f32_32x32x8f16_imm_splat(<16 x float> addrspace(1)* %arg) {
422bb:
423  %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x8f16(<4 x half> <half 1.0, half 1.0, half 1.0, half 1.0>, <4 x half> <half 2.0, half 2.0, half 2.0, half 2.0>, <16 x float> <float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0>, i32 0, i32 0, i32 0)
424  store <16 x float> %mai.1, <16 x float> addrspace(1)* %arg
425  ret void
426}
427
428; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_imm_splat:
429; GCN-DAG:         v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
430; GCN-DAG:         v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
431; NOLIT-SRCC-DAG:  v_accvgpr_write_b32 a{{[0-9]+}}, 0
432; NOLIT-SRCC:      v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9:]+}}]
433; LIT-SRCC:        v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], 0
434; GFX90A:          v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], 0
435; GFX908-COUNT-32: v_accvgpr_read_b32
436; GFX908:          global_store_dwordx4
437; GFX90A-NOT:      v_accvgpr_read_b32
438; GFX90A-COUNT-8:  global_store_dwordx4 {{v[0-9]+}}, a[{{[0-9:]+}}],
439define amdgpu_kernel void @test_mfma_f32_32x32x1f32_imm_splat(<32 x float> addrspace(1)* %arg) {
440bb:
441  %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> <float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0>, i32 0, i32 0, i32 0)
442  store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
443  ret void
444}
445
446; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32_imm:
447; GCN-DAG:        v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
448; GCN-DAG:        v_accvgpr_write_b32 a{{[0-9]+}}, 2.0
449; GFX908-DAG:     v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
450; GFX908-DAG:     v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
451; GFX90A-DAG:     v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
452; GFX90A-DAG:     v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
453; GFX908_A:       v_mfma_f32_4x4x1f32 [[RES:a\[[0-9]+:[0-9]+\]]], {{v[0-9]+}}, {{v[0-9]+}}, a[{{[0-9]+:[0-9]+}}]
454; GFX908-COUNT-4: v_accvgpr_read_b32
455; GFX908:         global_store_dwordx4
456; GFX90A-NOT:     v_accvgpr_read_b32
457; GFX90A:         global_store_dwordx4 {{v[0-9]+}}, [[RES]],
458define amdgpu_kernel void @test_mfma_f32_4x4x1f32_imm(<4 x float> addrspace(1)* %arg) {
459bb:
460  %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> <float 1.0, float 2.0, float 1.0, float 1.0>, i32 0, i32 0, i32 0)
461  store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
462  ret void
463}
464
465; GCN-LABEL: {{^}}test_mfma_f32_16x16x1f32_imm:
466; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
467; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, 2.0
468; GFX908-COUNT-14: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
469; GFX90A-COUNT-14: v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
470; GFX908_A:        v_mfma_f32_16x16x1f32 a[{{[0-9]+:[0-9]+}}], {{v[0-9]+}}, {{v[0-9]+}}, a[{{[0-9]+:[0-9]+}}]
471; GFX908-COUNT-16: v_accvgpr_read_b32
472; GFX908:          global_store_dwordx4
473; GFX90A-NOT:      v_accvgpr_read_b32
474; GFX90A-COUNT-4:  global_store_dwordx4 {{v[0-9]+}}, a[{{[0-9:]+}}],
475define amdgpu_kernel void @test_mfma_f32_16x16x1f32_imm(<16 x float> addrspace(1)* %arg) {
476bb:
477  %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x1f32(float 1.0, float 2.0, <16 x float> <float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 1.0, float 2.0>, i32 0, i32 0, i32 0)
478  store <16 x float> %mai.1, <16 x float> addrspace(1)* %arg
479  ret void
480}
481
482; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_imm:
483; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, 0
484; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, 1.0
485; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
486; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
487; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
488; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
489; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
490; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
491; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
492; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
493; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
494; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
495; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
496; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
497; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
498; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
499; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
500; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
501; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
502; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
503; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
504; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
505; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
506; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
507; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
508; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
509; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
510; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
511; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
512; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
513; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
514; GFX908-DAG:      v_accvgpr_write_b32 a{{[0-9]+}}, 0
515; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
516; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
517; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
518; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
519; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
520; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
521; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
522; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
523; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
524; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
525; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
526; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
527; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
528; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
529; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
530; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
531; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
532; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
533; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
534; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
535; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
536; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
537; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
538; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
539; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
540; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
541; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
542; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
543; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
544; GFX90A-DAG:      v_accvgpr_mov_b32 a{{[0-9]+}}, a{{[0-9]+}}
545; GFX908_A:        v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], {{v[0-9]+}}, {{v[0-9]+}}, a[{{[0-9]+:[0-9]+}}]
546; GFX908-COUNT-32: v_accvgpr_read_b32
547; GFX908:          global_store_dwordx4
548; GFX90A-NOT:      v_accvgpr_read_b32
549; GFX90A-COUNT-8:  global_store_dwordx4 {{v[0-9]+}}, a[{{[0-9:]+}}],
550define amdgpu_kernel void @test_mfma_f32_32x32x1f32_imm(<32 x float> addrspace(1)* %arg) {
551bb:
552  %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> <float 1.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0, float 0.0>, i32 0, i32 0, i32 0)
553  store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
554  ret void
555}
556
557; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32_lit_splat:
558; GFX908_A:       v_mov_b32_e32 [[TMP:v[0-9]+]], 0x42f60000
559; GCN:            v_accvgpr_write_b32 [[TTMPA:a[0-9]+]], [[TMP]]
560; GFX908:         v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]
561; GFX908:         v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]
562; GFX908:         v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]
563; GFX90A:         v_accvgpr_mov_b32 a{{[0-9]+}}, [[TTMPA]]
564; GFX90A:         v_accvgpr_mov_b32 a{{[0-9]+}}, [[TTMPA]]
565; GFX90A:         v_accvgpr_mov_b32 a{{[0-9]+}}, [[TTMPA]]
566; GFX908_A:       v_mfma_f32_4x4x1f32 [[RES:a\[[0-9]+:[0-9]+\]]], {{v[0-9]+}}, {{v[0-9]+}}, a[{{[0-9]+:[0-9]+}}]
567; GFX908-COUNT-4: v_accvgpr_read_b32
568; GFX908:         global_store_dwordx4
569; GFX90A-NOT:     v_accvgpr_read_b32
570; GFX90A:         global_store_dwordx4 {{v[0-9]+}}, [[RES]]
571define amdgpu_kernel void @test_mfma_f32_4x4x1f32_lit_splat(<4 x float> addrspace(1)* %arg, i64 %idx) {
572bb:
573  %tid = call i32 @llvm.amdgcn.workitem.id.x()
574  %gep = getelementptr inbounds <4 x float>, <4 x float> addrspace(1)* %arg, i32 %tid
575  %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> <float 123.0, float 123.0, float 123.0, float 123.0>, i32 0, i32 0, i32 0)
576  ;store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
577  store <4 x float> %mai.1, <4 x float> addrspace(1)* %gep
578  ret void
579}
580
581; FIXME: Resulting code for splat is pretty bad. A v_mov_b32 is moved
582; in the middle of the expanded agpr reg_sequence. The broadcast of
583; the individual AGPR->AGPR components should avoid the intermediate AGPR case.
584; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32_lit_splat_bad_code:
585; GFX908_A: v_mov_b32_e32 [[TMP0:v[0-9]+]], 0x42f60000
586; GCN:      v_accvgpr_write_b32 [[AGPR:a[0-9]+]], [[TMP0]]
587; GFX908:   s_nop 0
588; GFX908:   v_accvgpr_read_b32 [[TMP1:v[0-9]+]], [[AGPR]]
589; GFX908:   v_accvgpr_read_b32 [[TMP2:v[0-9]+]], [[AGPR]]
590; GFX908:   v_accvgpr_read_b32 [[TMP3:v[0-9]+]], [[AGPR]]
591; GFX908:   v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
592; GFX908:   v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
593; GFX908:   v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
594; GFX90A-COUNT-3: v_accvgpr_mov_b32 a{{[0-9]+}}, [[AGPR]]
595; GCN: s_nop 0
596; GFX908_A:  v_mfma_f32_4x4x1f32 a[{{[0-9]+:[0-9]+}}], {{v[0-9]+}}, {{v[0-9]+}}, a[{{[0-9]+:[0-9]+}}]
597; GFX908-COUNT-4: v_accvgpr_read_b32
598; GFX908:    global_store_dwordx4 v{{[0-9]+}}, v[{{[0-9:]+}}], s[{{[0-9:]+}}]
599; GFX90A:    global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}], s[{{[0-9:]+}}]
600define amdgpu_kernel void @test_mfma_f32_4x4x1f32_lit_splat_bad_code(<4 x float> addrspace(1)* %arg) {
601bb:
602  %tid = call i32 @llvm.amdgcn.workitem.id.x()
603  %gep = getelementptr inbounds <4 x float>, <4 x float> addrspace(1)* %arg, i32 %tid
604
605  %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> <float 123.0, float 123.0, float 123.0, float 123.0>, i32 0, i32 0, i32 0)
606  store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
607  ret void
608}
609
610; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_vecarg:
611; GFX90A-DAG:      v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
612; GCN-DAG:         v_mov_b32_e32 [[ONE:v[0-9]+]], 1.0
613; GCN-COUNT-8:     global_load_dwordx4
614; GFX908-COUNT-16: v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
615; GFX908-DAG:      v_mov_b32_e32 [[TWO:v[0-9]+]], 2.0
616; GFX90A-NOT:      v_accvgpr_write
617; GFX908:          v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
618; GFX90A:          v_mfma_f32_32x32x1f32 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
619; GFX908-COUNT-32: v_accvgpr_read_b32
620; GFX908-COUNT-8:  global_store_dwordx4
621; GFX90A-NOT:      v_accvgpr_read_b32
622; GFX90A-COUNT-5:  global_store_dwordx4 v{{[0-9:]+}}, a[{{[0-9:]+}}], s[{{[0-9:]+}}]
623define amdgpu_kernel void @test_mfma_f32_32x32x1f32_vecarg(<32 x float> addrspace(1)* %arg) {
624bb:
625  %tid = call i32 @llvm.amdgcn.workitem.id.x()
626  %gep = getelementptr inbounds <32 x float>, <32 x float> addrspace(1)* %arg, i32 %tid
627  %in.1 = load <32 x float>, <32 x float> addrspace(1)* %gep
628  %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> %in.1, i32 1, i32 2, i32 3)
629  store <32 x float> %mai.1, <32 x float> addrspace(1)* %gep
630  ret void
631}
632