xref: /llvm-project/llvm/test/CodeGen/AMDGPU/attr-amdgpu-flat-work-group-size.ll (revision 4490003a22658dcd12527029b2c8682b63d8a9d6) 
1; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -verify-machineinstrs < %s | FileCheck --check-prefix=CHECK %s
2; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -verify-machineinstrs -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefix=HSAMD %s
3
4; CHECK-LABEL: {{^}}min_64_max_64:
5; CHECK: SGPRBlocks: 0
6; CHECK: VGPRBlocks: 0
7; CHECK: NumSGPRsForWavesPerEU: 1
8; CHECK: NumVGPRsForWavesPerEU: 1
9define amdgpu_kernel void @min_64_max_64() #0 {
10entry:
11  ret void
12}
13attributes #0 = {"amdgpu-flat-work-group-size"="64,64"}
14
15; CHECK-LABEL: {{^}}min_64_max_128:
16; CHECK: SGPRBlocks: 0
17; CHECK: VGPRBlocks: 0
18; CHECK: NumSGPRsForWavesPerEU: 1
19; CHECK: NumVGPRsForWavesPerEU: 1
20define amdgpu_kernel void @min_64_max_128() #1 {
21entry:
22  ret void
23}
24attributes #1 = {"amdgpu-flat-work-group-size"="64,128"}
25
26; CHECK-LABEL: {{^}}min_128_max_128:
27; CHECK: SGPRBlocks: 0
28; CHECK: VGPRBlocks: 0
29; CHECK: NumSGPRsForWavesPerEU: 1
30; CHECK: NumVGPRsForWavesPerEU: 1
31define amdgpu_kernel void @min_128_max_128() #2 {
32entry:
33  ret void
34}
35attributes #2 = {"amdgpu-flat-work-group-size"="128,128"}
36
37; CHECK-LABEL: {{^}}min_1024_max_1024
38; CHECK: SGPRBlocks: 0
39; CHECK: VGPRBlocks: 10
40; CHECK: NumSGPRsForWavesPerEU: 2{{$}}
41; CHECK: NumVGPRsForWavesPerEU: 43
42@var = addrspace(1) global float 0.0
43define amdgpu_kernel void @min_1024_max_1024() #3 {
44  %val0 = load volatile float, ptr addrspace(1) @var
45  %val1 = load volatile float, ptr addrspace(1) @var
46  %val2 = load volatile float, ptr addrspace(1) @var
47  %val3 = load volatile float, ptr addrspace(1) @var
48  %val4 = load volatile float, ptr addrspace(1) @var
49  %val5 = load volatile float, ptr addrspace(1) @var
50  %val6 = load volatile float, ptr addrspace(1) @var
51  %val7 = load volatile float, ptr addrspace(1) @var
52  %val8 = load volatile float, ptr addrspace(1) @var
53  %val9 = load volatile float, ptr addrspace(1) @var
54  %val10 = load volatile float, ptr addrspace(1) @var
55  %val11 = load volatile float, ptr addrspace(1) @var
56  %val12 = load volatile float, ptr addrspace(1) @var
57  %val13 = load volatile float, ptr addrspace(1) @var
58  %val14 = load volatile float, ptr addrspace(1) @var
59  %val15 = load volatile float, ptr addrspace(1) @var
60  %val16 = load volatile float, ptr addrspace(1) @var
61  %val17 = load volatile float, ptr addrspace(1) @var
62  %val18 = load volatile float, ptr addrspace(1) @var
63  %val19 = load volatile float, ptr addrspace(1) @var
64  %val20 = load volatile float, ptr addrspace(1) @var
65  %val21 = load volatile float, ptr addrspace(1) @var
66  %val22 = load volatile float, ptr addrspace(1) @var
67  %val23 = load volatile float, ptr addrspace(1) @var
68  %val24 = load volatile float, ptr addrspace(1) @var
69  %val25 = load volatile float, ptr addrspace(1) @var
70  %val26 = load volatile float, ptr addrspace(1) @var
71  %val27 = load volatile float, ptr addrspace(1) @var
72  %val28 = load volatile float, ptr addrspace(1) @var
73  %val29 = load volatile float, ptr addrspace(1) @var
74  %val30 = load volatile float, ptr addrspace(1) @var
75  %val31 = load volatile float, ptr addrspace(1) @var
76  %val32 = load volatile float, ptr addrspace(1) @var
77  %val33 = load volatile float, ptr addrspace(1) @var
78  %val34 = load volatile float, ptr addrspace(1) @var
79  %val35 = load volatile float, ptr addrspace(1) @var
80  %val36 = load volatile float, ptr addrspace(1) @var
81  %val37 = load volatile float, ptr addrspace(1) @var
82  %val38 = load volatile float, ptr addrspace(1) @var
83  %val39 = load volatile float, ptr addrspace(1) @var
84  %val40 = load volatile float, ptr addrspace(1) @var
85
86  store volatile float %val0, ptr addrspace(1) @var
87  store volatile float %val1, ptr addrspace(1) @var
88  store volatile float %val2, ptr addrspace(1) @var
89  store volatile float %val3, ptr addrspace(1) @var
90  store volatile float %val4, ptr addrspace(1) @var
91  store volatile float %val5, ptr addrspace(1) @var
92  store volatile float %val6, ptr addrspace(1) @var
93  store volatile float %val7, ptr addrspace(1) @var
94  store volatile float %val8, ptr addrspace(1) @var
95  store volatile float %val9, ptr addrspace(1) @var
96  store volatile float %val10, ptr addrspace(1) @var
97  store volatile float %val11, ptr addrspace(1) @var
98  store volatile float %val12, ptr addrspace(1) @var
99  store volatile float %val13, ptr addrspace(1) @var
100  store volatile float %val14, ptr addrspace(1) @var
101  store volatile float %val15, ptr addrspace(1) @var
102  store volatile float %val16, ptr addrspace(1) @var
103  store volatile float %val17, ptr addrspace(1) @var
104  store volatile float %val18, ptr addrspace(1) @var
105  store volatile float %val19, ptr addrspace(1) @var
106  store volatile float %val20, ptr addrspace(1) @var
107  store volatile float %val21, ptr addrspace(1) @var
108  store volatile float %val22, ptr addrspace(1) @var
109  store volatile float %val23, ptr addrspace(1) @var
110  store volatile float %val24, ptr addrspace(1) @var
111  store volatile float %val25, ptr addrspace(1) @var
112  store volatile float %val26, ptr addrspace(1) @var
113  store volatile float %val27, ptr addrspace(1) @var
114  store volatile float %val28, ptr addrspace(1) @var
115  store volatile float %val29, ptr addrspace(1) @var
116  store volatile float %val30, ptr addrspace(1) @var
117  store volatile float %val31, ptr addrspace(1) @var
118  store volatile float %val32, ptr addrspace(1) @var
119  store volatile float %val33, ptr addrspace(1) @var
120  store volatile float %val34, ptr addrspace(1) @var
121  store volatile float %val35, ptr addrspace(1) @var
122  store volatile float %val36, ptr addrspace(1) @var
123  store volatile float %val37, ptr addrspace(1) @var
124  store volatile float %val38, ptr addrspace(1) @var
125  store volatile float %val39, ptr addrspace(1) @var
126  store volatile float %val40, ptr addrspace(1) @var
127
128  ret void
129}
130attributes #3 = {"amdgpu-flat-work-group-size"="1024,1024"}
131
132!llvm.module.flags = !{!0}
133!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}
134
135; HSAMD: amdhsa.kernels
136; HSAMD:  .max_flat_workgroup_size: 64
137; HSAMD:  .name: min_64_max_64
138; HSAMD:  .max_flat_workgroup_size: 128
139; HSAMD:  .name: min_64_max_128
140; HSAMD:  .max_flat_workgroup_size: 128
141; HSAMD:  .name: min_128_max_128
142; HSAMD:  .max_flat_workgroup_size: 1024
143; HSAMD:  .name: min_1024_max_1024
144