1// RUN: %check_clang_tidy -check-suffix=OLDCLOLDAOC %s altera-single-work-item-barrier %t -- -header-filter=.* "--" -cl-std=CL1.2 -c -DOLDCLOLDAOC
2// RUN: %check_clang_tidy -check-suffix=NEWCLOLDAOC %s altera-single-work-item-barrier %t -- -header-filter=.* "--" -cl-std=CL2.0 -c -DNEWCLOLDAOC
3// RUN: %check_clang_tidy -check-suffix=OLDCLNEWAOC %s altera-single-work-item-barrier %t -- -config='{CheckOptions: { altera-single-work-item-barrier.AOCVersion: 1701}}' -header-filter=.* "--" -cl-std=CL1.2 -c -DOLDCLNEWAOC
4// RUN: %check_clang_tidy -check-suffix=NEWCLNEWAOC %s altera-single-work-item-barrier %t -- -config='{CheckOptions: { altera-single-work-item-barrier.AOCVersion: 1701}}' -header-filter=.* "--" -cl-std=CL2.0 -c -DNEWCLNEWAOC
5
6#ifdef OLDCLOLDAOC // OpenCL 1.2 Altera Offline Compiler < 17.1
7void __kernel error_barrier_no_id(__global int * foo, int size) {
8 // CHECK-MESSAGES-OLDCLOLDAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id' does not call 'get_global_id' or 'get_local_id' and will be treated as a single work-item [altera-single-work-item-barrier]
9 for (int j = 0; j < 256; j++) {
10 for (int i = 256; i < size; i+= 256) {
11 foo[j] += foo[j+i];
12 }
13 }
14 barrier(CLK_GLOBAL_MEM_FENCE);
15 // CHECK-MESSAGES-OLDCLOLDAOC: :[[@LINE-1]]:3: note: barrier call is in a single work-item and may error out
16 for (int i = 1; i < 256; i++) {
17 foo[0] += foo[i];
18 }
19}
20
21void __kernel success_barrier_global_id(__global int * foo, int size) {
22 barrier(CLK_GLOBAL_MEM_FENCE);
23 int tid = get_global_id(0);
24}
25
26void __kernel success_barrier_local_id(__global int * foo, int size) {
27 barrier(CLK_GLOBAL_MEM_FENCE);
28 int tid = get_local_id(0);
29}
30
31void __kernel success_barrier_both_ids(__global int * foo, int size) {
32 barrier(CLK_GLOBAL_MEM_FENCE);
33 int gid = get_global_id(0);
34 int lid = get_local_id(0);
35}
36
37void success_nokernel_barrier_no_id(__global int * foo, int size) {
38 for (int j = 0; j < 256; j++) {
39 for (int i = 256; i < size; i+= 256) {
40 foo[j] += foo[j+i];
41 }
42 }
43 barrier(CLK_GLOBAL_MEM_FENCE);
44 for (int i = 1; i < 256; i++) {
45 foo[0] += foo[i];
46 }
47}
48
49void success_nokernel_barrier_global_id(__global int * foo, int size) {
50 barrier(CLK_GLOBAL_MEM_FENCE);
51 int tid = get_global_id(0);
52}
53
54void success_nokernel_barrier_local_id(__global int * foo, int size) {
55 barrier(CLK_GLOBAL_MEM_FENCE);
56 int tid = get_local_id(0);
57}
58
59void success_nokernel_barrier_both_ids(__global int * foo, int size) {
60 barrier(CLK_GLOBAL_MEM_FENCE);
61 int gid = get_global_id(0);
62 int lid = get_local_id(0);
63}
64#endif
65
66#ifdef NEWCLOLDAOC // OpenCL 2.0 Altera Offline Compiler < 17.1
67void __kernel error_barrier_no_id(__global int * foo, int size) {
68 // CHECK-MESSAGES-NEWCLOLDAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id' does not call 'get_global_id' or 'get_local_id' and will be treated as a single work-item [altera-single-work-item-barrier]
69 for (int j = 0; j < 256; j++) {
70 for (int i = 256; i < size; i+= 256) {
71 foo[j] += foo[j+i];
72 }
73 }
74 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
75 // CHECK-MESSAGES-NEWCLOLDAOC: :[[@LINE-1]]:3: note: barrier call is in a single work-item and may error out
76 for (int i = 1; i < 256; i++) {
77 foo[0] += foo[i];
78 }
79}
80
81void __kernel success_barrier_global_id(__global int * foo, int size) {
82 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
83 int tid = get_global_id(0);
84}
85
86void __kernel success_barrier_local_id(__global int * foo, int size) {
87 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
88 int tid = get_local_id(0);
89}
90
91void __kernel success_barrier_both_ids(__global int * foo, int size) {
92 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
93 int gid = get_global_id(0);
94 int lid = get_local_id(0);
95}
96
97void success_nokernel_barrier_no_id(__global int * foo, int size) {
98 for (int j = 0; j < 256; j++) {
99 for (int i = 256; i < size; i+= 256) {
100 foo[j] += foo[j+i];
101 }
102 }
103 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
104 for (int i = 1; i < 256; i++) {
105 foo[0] += foo[i];
106 }
107}
108
109void success_nokernel_barrier_global_id(__global int * foo, int size) {
110 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
111 int tid = get_global_id(0);
112}
113
114void success_nokernel_barrier_local_id(__global int * foo, int size) {
115 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
116 int tid = get_local_id(0);
117}
118
119void success_nokernel_barrier_both_ids(__global int * foo, int size) {
120 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
121 int gid = get_global_id(0);
122 int lid = get_local_id(0);
123}
124#endif
125
126#ifdef OLDCLNEWAOC // OpenCL 1.2 Altera Offline Compiler >= 17.1
127void __kernel error_barrier_no_id(__global int * foo, int size) {
128 // CHECK-MESSAGES-OLDCLNEWAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id' does not call an ID function and may be a viable single work-item, but will be forced to execute as an NDRange [altera-single-work-item-barrier]
129 for (int j = 0; j < 256; j++) {
130 for (int i = 256; i < size; i+= 256) {
131 foo[j] += foo[j+i];
132 }
133 }
134 barrier(CLK_GLOBAL_MEM_FENCE);
135 // CHECK-MESSAGES-OLDCLNEWAOC: :[[@LINE-1]]:3: note: barrier call will force NDRange execution; if single work-item semantics are desired a mem_fence may be more efficient
136 for (int i = 1; i < 256; i++) {
137 foo[0] += foo[i];
138 }
139}
140
141__attribute__ ((reqd_work_group_size(1,1,1)))
142void __kernel error_barrier_no_id_work_group_size(__global int * foo, int size) {
143 // CHECK-MESSAGES-OLDCLNEWAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id_work_group_size' does not call an ID function and may be a viable single work-item, but will be forced to execute as an NDRange [altera-single-work-item-barrier]
144 for (int j = 0; j < 256; j++) {
145 for (int i = 256; i < size; i+= 256) {
146 foo[j] += foo[j+i];
147 }
148 }
149 barrier(CLK_GLOBAL_MEM_FENCE);
150 // CHECK-MESSAGES-OLDCLNEWAOC: :[[@LINE-1]]:3: note: barrier call will force NDRange execution; if single work-item semantics are desired a mem_fence may be more efficient
151 for (int i = 1; i < 256; i++) {
152 foo[0] += foo[i];
153 }
154}
155
156__attribute__ ((reqd_work_group_size(2,1,1)))
157void __kernel success_barrier_no_id_work_group_size(__global int * foo, int size) {
158 for (int j = 0; j < 256; j++) {
159 for (int i = 256; i < size; i+= 256) {
160 foo[j] += foo[j+i];
161 }
162 }
163 barrier(CLK_GLOBAL_MEM_FENCE);
164 for (int i = 1; i < 256; i++) {
165 foo[0] += foo[i];
166 }
167}
168
169void __kernel success_barrier_global_id(__global int * foo, int size) {
170 barrier(CLK_GLOBAL_MEM_FENCE);
171 int tid = get_global_id(0);
172}
173
174void __kernel success_barrier_local_id(__global int * foo, int size) {
175 barrier(CLK_GLOBAL_MEM_FENCE);
176 int tid = get_local_id(0);
177}
178
179void __kernel success_barrier_both_ids(__global int * foo, int size) {
180 barrier(CLK_GLOBAL_MEM_FENCE);
181 int gid = get_global_id(0);
182 int lid = get_local_id(0);
183}
184
185void success_nokernel_barrier_no_id(__global int * foo, int size) {
186 for (int j = 0; j < 256; j++) {
187 for (int i = 256; i < size; i+= 256) {
188 foo[j] += foo[j+i];
189 }
190 }
191 barrier(CLK_GLOBAL_MEM_FENCE);
192 for (int i = 1; i < 256; i++) {
193 foo[0] += foo[i];
194 }
195}
196
197void success_nokernel_barrier_global_id(__global int * foo, int size) {
198 barrier(CLK_GLOBAL_MEM_FENCE);
199 int tid = get_global_id(0);
200}
201
202void success_nokernel_barrier_local_id(__global int * foo, int size) {
203 barrier(CLK_GLOBAL_MEM_FENCE);
204 int tid = get_local_id(0);
205}
206
207void success_nokernel_barrier_both_ids(__global int * foo, int size) {
208 barrier(CLK_GLOBAL_MEM_FENCE);
209 int gid = get_global_id(0);
210 int lid = get_local_id(0);
211}
212#endif
213
214#ifdef NEWCLNEWAOC // OpenCL 2.0 Altera Offline Compiler >= 17.1
215void __kernel error_barrier_no_id(__global int * foo, int size) {
216 // CHECK-MESSAGES-NEWCLNEWAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id' does not call an ID function and may be a viable single work-item, but will be forced to execute as an NDRange [altera-single-work-item-barrier]
217 for (int j = 0; j < 256; j++) {
218 for (int i = 256; i < size; i+= 256) {
219 foo[j] += foo[j+i];
220 }
221 }
222 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
223 // CHECK-MESSAGES-NEWCLNEWAOC: :[[@LINE-1]]:3: note: barrier call will force NDRange execution; if single work-item semantics are desired a mem_fence may be more efficient
224 for (int i = 1; i < 256; i++) {
225 foo[0] += foo[i];
226 }
227}
228
229__attribute__ ((reqd_work_group_size(1,1,1)))
230void __kernel error_barrier_no_id_work_group_size(__global int * foo, int size) {
231 // CHECK-MESSAGES-NEWCLNEWAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id_work_group_size' does not call an ID function and may be a viable single work-item, but will be forced to execute as an NDRange [altera-single-work-item-barrier]
232 for (int j = 0; j < 256; j++) {
233 for (int i = 256; i < size; i+= 256) {
234 foo[j] += foo[j+i];
235 }
236 }
237 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
238 // CHECK-MESSAGES-NEWCLNEWAOC: :[[@LINE-1]]:3: note: barrier call will force NDRange execution; if single work-item semantics are desired a mem_fence may be more efficient
239 for (int i = 1; i < 256; i++) {
240 foo[0] += foo[i];
241 }
242}
243
244__attribute__ ((reqd_work_group_size(2,1,1)))
245void __kernel success_barrier_no_id_work_group_size(__global int * foo, int size) {
246 for (int j = 0; j < 256; j++) {
247 for (int i = 256; i < size; i+= 256) {
248 foo[j] += foo[j+i];
249 }
250 }
251 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
252 for (int i = 1; i < 256; i++) {
253 foo[0] += foo[i];
254 }
255}
256
257void __kernel success_barrier_global_id(__global int * foo, int size) {
258 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
259 int tid = get_global_id(0);
260}
261
262void __kernel success_barrier_local_id(__global int * foo, int size) {
263 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
264 int tid = get_local_id(0);
265}
266
267void __kernel success_barrier_both_ids(__global int * foo, int size) {
268 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
269 int gid = get_global_id(0);
270 int lid = get_local_id(0);
271}
272
273void success_nokernel_barrier_no_id(__global int * foo, int size) {
274 for (int j = 0; j < 256; j++) {
275 for (int i = 256; i < size; i+= 256) {
276 foo[j] += foo[j+i];
277 }
278 }
279 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
280 for (int i = 1; i < 256; i++) {
281 foo[0] += foo[i];
282 }
283}
284
285void success_nokernel_barrier_global_id(__global int * foo, int size) {
286 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
287 int tid = get_global_id(0);
288}
289
290void success_nokernel_barrier_local_id(__global int * foo, int size) {
291 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
292 int tid = get_local_id(0);
293}
294
295void success_nokernel_barrier_both_ids(__global int * foo, int size) {
296 work_group_barrier(CLK_GLOBAL_MEM_FENCE);
297 int gid = get_global_id(0);
298 int lid = get_local_id(0);
299}
300#endif
301

source code of clang-tools-extra/test/clang-tidy/checkers/altera/single-work-item-barrier.cpp