1; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck --enable-var-scope --check-prefixes=GCN,GFX908 %s
2; RUN: llc -march=amdgcn -mcpu=gfx90a -verify-machineinstrs < %s | FileCheck --enable-var-scope --check-prefixes=GCN,GFX90A %s
3; RUN: llc -global-isel -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.32x32x1f32(float, float, <32 x float>, i32, i32, i32)
6
7; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_vgpr:
8; GFX908: v_mfma_f32_32x32x1{{.*}} a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9:]+}}, a[{{[0-9:]+}}]
9; GFX90A: v_mfma_f32_32x32x1{{.*}} v[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9:]+}}, v[{{[0-9:]+}}]
10define amdgpu_kernel void @test_mfma_f32_32x32x1f32_vgpr(<32 x float> addrspace(1)* %arg) #0 {
11bb:
12  %in.1 = load <32 x float>, <32 x float> addrspace(1)* %arg
13  %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)
14  store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
15  ret void
16}
17
18; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_agpr:
19; GCN: v_mfma_f32_32x32x1{{.*}} a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9:]+}}, a[{{[0-9:]+}}]
20define amdgpu_kernel void @test_mfma_f32_32x32x1f32_agpr(<32 x float> addrspace(1)* %arg) #1 {
21bb:
22  %in.1 = load <32 x float>, <32 x float> addrspace(1)* %arg
23  %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)
24  store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
25  ret void
26}
27
28; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_inline_asm_virtual_agpr:
29; GCN: v_mfma_f32_32x32x1{{.*}} a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9:]+}}, a[{{[0-9:]+}}]
30define amdgpu_kernel void @test_mfma_f32_32x32x1f32_inline_asm_virtual_agpr(<32 x float> addrspace(1)* %arg) #0 {
31bb:
32  %acc = call i32 asm sideeffect "; def $0", "={a0}"()
33  %in.1 = load <32 x float>, <32 x float> addrspace(1)* %arg
34  %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)
35  store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
36  ret void
37}
38
39; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_inline_asm_phys_agpr:
40; GCN: v_mfma_f32_32x32x1{{.*}} a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9:]+}}, a[{{[0-9:]+}}]
41define amdgpu_kernel void @test_mfma_f32_32x32x1f32_inline_asm_phys_agpr(<32 x float> addrspace(1)* %arg) #0 {
42bb:
43  call void asm sideeffect "; use $0", "{a[100:131]}"(<32 x float> undef)
44  %in.1 = load <32 x float>, <32 x float> addrspace(1)* %arg
45  %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)
46  store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
47  ret void
48}
49
50; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_inline_asm_no_agprs:
51; GFX908: v_mfma_f32_32x32x1{{.*}} a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9:]+}}, a[{{[0-9:]+}}]
52; GFX90A: v_mfma_f32_32x32x1{{.*}} v[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9:]+}}, v[{{[0-9:]+}}]
53define amdgpu_kernel void @test_mfma_f32_32x32x1f32_inline_asm_no_agprs(<32 x float> addrspace(1)* %arg) #0 {
54bb:
55  %acc = call i32 asm sideeffect "; def $0", "={v0}"()
56  %in.1 = load <32 x float>, <32 x float> addrspace(1)* %arg
57  %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)
58  store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
59  ret void
60}
61
62; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_call:
63; GCN: v_mfma_f32_32x32x1{{.*}} a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9:]+}}, a[{{[0-9:]+}}]
64define amdgpu_kernel void @test_mfma_f32_32x32x1f32_call(<32 x float> addrspace(1)* %arg) #0 {
65bb:
66  call void @foo()
67  %in.1 = load <32 x float>, <32 x float> addrspace(1)* %arg
68  %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)
69  store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
70  ret void
71}
72
73; We could avoid scan to find calls since we see these during lowering before selection.
74; However, in SDag lowering and selection is done block by block, so it would only work
75; in Global ISel.
76
77; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_call_multi_bb:
78; GCN: v_mfma_f32_32x32x1{{.*}} a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9:]+}}, a[{{[0-9:]+}}]
79define amdgpu_kernel void @test_mfma_f32_32x32x1f32_call_multi_bb(<32 x float> addrspace(1)* %arg) #0 {
80bb1:
81  %in.1 = load <32 x float>, <32 x float> addrspace(1)* %arg
82  %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)
83  store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
84  br i1 undef, label %bb2, label %bb3
85  br label %bb2
86
87bb2:
88  call void @foo()
89  br label %bb3
90
91bb3:
92  ret void
93}
94
95; GCN-LABEL: {{^}}test_mfma_f32_32x32x1f32_nonentry:
96; GCN: v_mfma_f32_32x32x1{{.*}} a[{{[0-9:]+}}], v{{[0-9]+}}, v{{[0-9:]+}}, a[{{[0-9:]+}}]
97define void @test_mfma_f32_32x32x1f32_nonentry(<32 x float> addrspace(1)* %arg) #0 {
98bb:
99  %in.1 = load <32 x float>, <32 x float> addrspace(1)* %arg
100  %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)
101  store <32 x float> %mai.1, <32 x float> addrspace(1)* %arg
102  ret void
103}
104
105declare void @foo()
106
107attributes #0 = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-waves-per-eu"="2" }
108attributes #1 = { "amdgpu-flat-work-group-size"="1,256" }
109