1; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN %s
2; RUN: llc -march=amdgcn -mcpu=gfx908 -mattr=-mfma-inline-literal-bug -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN %s
3; RUN: llc -march=amdgcn -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX90A %s
4
5declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x2bf16(<2 x i16>, <2 x i16>, <32 x float>, i32, i32, i32)
6declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x2bf16(<2 x i16>, <2 x i16>, <16 x float>, i32, i32, i32)
7declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16>, <2 x i16>, <4 x float>, i32, i32, i32)
8declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16(<2 x i16>, <2 x i16>, <16 x float>, i32, i32, i32)
9declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x8bf16(<2 x i16>, <2 x i16>, <4 x float>, i32, i32, i32)
10declare i32 @llvm.amdgcn.workitem.id.x()
11
12; GCN-LABEL: {{^}}test_mfma_f32_32x32x2bf16:
13; GCN-DAG:         v_mov_b32_e32 [[TWO:v[0-9]+]], 2
14; GCN-DAG:         v_mov_b32_e32 [[ONE:v[0-9]+]], 1
15; GCN-DAG:         s_load_dwordx16
16; GCN-DAG:         s_load_dwordx16
17; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
18; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
19; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
20; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
21; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
22; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
23; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
24; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
25; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
26; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
27; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
28; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
29; GCN-DAG:         v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
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:             v_mfma_f32_32x32x2bf16 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
50; GFX908-COUNT-32: v_accvgpr_read_b32
51; GFX908:          global_store_dwordx4
52; GFX90A-NOT:      v_accvgpr_read_b32
53; GFX90A-COUNT-8:  global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}],
54define amdgpu_kernel void @test_mfma_f32_32x32x2bf16(<32 x float> addrspace(1)* %arg) {
55bb:
56  %in.1 = load <32 x float>, <32 x float> addrspace(1)* %arg
57  %a = bitcast i32 1 to <2 x i16>
58  %b = bitcast i32 2 to <2 x i16>
59  %mai.1 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x2bf16(<2 x i16> %a, <2 x i16> %b, <32 x float> %in.1, i32 1, i32 2, i32 3)
60  store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
61  ret void
62}
63
64; GCN-LABEL: {{^}}test_mfma_f32_16x16x2bf16:
65; GCN-DAG:         v_mov_b32_e32 [[TWO:v[0-9]+]], 2
66; GCN-DAG:         v_mov_b32_e32 [[ONE:v[0-9]+]], 1
67; GCN:             s_load_dwordx16
68; GCN-COUNT-16:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
69; GCN:             v_mfma_f32_16x16x2bf16 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
70; GFX908-COUNT-16: v_accvgpr_read_b32
71; GFX908:          global_store_dwordx4
72; GFX90A-NOT:      v_accvgpr_read_b32
73; GFX90A-COUNT-4:  global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}],
74define amdgpu_kernel void @test_mfma_f32_16x16x2bf16(<16 x float> addrspace(1)* %arg) {
75bb:
76  %in.1 = load <16 x float>, <16 x float> addrspace(1)* %arg
77  %a = bitcast i32 1 to <2 x i16>
78  %b = bitcast i32 2 to <2 x i16>
79  %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.16x16x2bf16(<2 x i16> %a, <2 x i16> %b, <16 x float> %in.1, i32 1, i32 2, i32 3)
80  store <16 x float> %mai.1, <16 x float> addrspace(1)* %arg
81  ret void
82}
83
84; GCN-LABEL: {{^}}test_mfma_f32_4x4x2bf16:
85; GCN-DAG:        v_mov_b32_e32 [[TWO:v[0-9]+]], 2
86; GCN-DAG:        v_mov_b32_e32 [[ONE:v[0-9]+]], 1
87; GCN:            s_load_dwordx4
88; GCN-COUNT-4:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
89; GCN:            v_mfma_f32_4x4x2bf16 [[RES:a\[[0-9]+:[0-9]+\]]], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
90; GFX908-COUNT-4: v_accvgpr_read_b32
91; GFX908:         global_store_dwordx4
92; GFX90A-NOT:     v_accvgpr_read_b32
93; GFX90A:         global_store_dwordx4 v{{[0-9]+}}, [[RES]],
94define amdgpu_kernel void @test_mfma_f32_4x4x2bf16(<4 x float> addrspace(1)* %arg) {
95bb:
96  %in.1 = load <4 x float>, <4 x float> addrspace(1)* %arg
97  %a = bitcast i32 1 to <2 x i16>
98  %b = bitcast i32 2 to <2 x i16>
99  %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x2bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %in.1, i32 1, i32 2, i32 3)
100  store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
101  ret void
102}
103
104; GCN-LABEL: {{^}}test_mfma_f32_32x32x4bf16:
105; GCN-DAG:         v_mov_b32_e32 [[TWO:v[0-9]+]], 2
106; GCN-DAG:         v_mov_b32_e32 [[ONE:v[0-9]+]], 1
107; GCN:             s_load_dwordx16
108; GCN-COUNT-16:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
109; GCN:             v_mfma_f32_32x32x4bf16 a[{{[0-9]+:[0-9]+}}], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
110; GFX908-COUNT-16: v_accvgpr_read_b32
111; GFX908:          global_store_dwordx4
112; GFX90A-NOT:      v_accvgpr_read_b32
113; GFX90A-COUNT-4:  global_store_dwordx4 v{{[0-9]+}}, a[{{[0-9:]+}}],
114define amdgpu_kernel void @test_mfma_f32_32x32x4bf16(<16 x float> addrspace(1)* %arg) {
115bb:
116  %in.1 = load <16 x float>, <16 x float> addrspace(1)* %arg
117  %a = bitcast i32 1 to <2 x i16>
118  %b = bitcast i32 2 to <2 x i16>
119  %mai.1 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16(<2 x i16> %a, <2 x i16> %b, <16 x float> %in.1, i32 1, i32 2, i32 3)
120  store <16 x float> %mai.1, <16 x float> addrspace(1)* %arg
121  ret void
122}
123
124; GCN-LABEL: {{^}}test_mfma_f32_16x16x8bf16:
125; GCN-DAG:        v_mov_b32_e32 [[TWO:v[0-9]+]], 2
126; GCN-DAG:        v_mov_b32_e32 [[ONE:v[0-9]+]], 1
127; GCN:            s_load_dwordx4
128; GCN-COUNT-4:    v_accvgpr_write_b32 a{{[0-9]+}}, v{{[0-9]+}}
129; GCN:            v_mfma_f32_16x16x8bf16 [[RES:a\[[0-9]+:[0-9]+\]]], [[ONE]], [[TWO]], a[{{[0-9]+:[0-9]+}}] cbsz:1 abid:2 blgp:3
130; GFX908-COUNT-4: v_accvgpr_read_b32
131; GFX908:         global_store_dwordx4
132; GFX90A-NOT:     v_accvgpr_read_b32
133; GFX90A:         global_store_dwordx4 v{{[0-9]+}}, [[RES]],
134define amdgpu_kernel void @test_mfma_f32_16x16x8bf16(<4 x float> addrspace(1)* %arg) {
135bb:
136  %in.1 = load <4 x float>, <4 x float> addrspace(1)* %arg
137  %a = bitcast i32 1 to <2 x i16>
138  %b = bitcast i32 2 to <2 x i16>
139  %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x8bf16(<2 x i16> %a, <2 x i16> %b, <4 x float> %in.1, i32 1, i32 2, i32 3)
140  store <4 x float> %mai.1, <4 x float> addrspace(1)* %arg
141  ret void
142}
143