1 // RUN: %check_clang_tidy -check-suffix=OLDCLOLDAOC %s altera-single-work-item-barrier %t -- -header-filter=.* "--" -cl-std=CL1.2 -c --include opencl-c.h -DOLDCLOLDAOC
2 // RUN: %check_clang_tidy -check-suffix=NEWCLOLDAOC %s altera-single-work-item-barrier %t -- -header-filter=.* "--" -cl-std=CL2.0 -c --include opencl-c.h -DNEWCLOLDAOC
3 // RUN: %check_clang_tidy -check-suffix=OLDCLNEWAOC %s altera-single-work-item-barrier %t -- -config='{CheckOptions: [{key: altera-single-work-item-barrier.AOCVersion, value: 1701}]}' -header-filter=.* "--" -cl-std=CL1.2 -c --include opencl-c.h -DOLDCLNEWAOC
4 // RUN: %check_clang_tidy -check-suffix=NEWCLNEWAOC %s altera-single-work-item-barrier %t -- -config='{CheckOptions: [{key: altera-single-work-item-barrier.AOCVersion, value: 1701}]}' -header-filter=.* "--" -cl-std=CL2.0 -c --include opencl-c.h -DNEWCLNEWAOC
5
6 #ifdef OLDCLOLDAOC // OpenCL 1.2 Altera Offline Compiler < 17.1
error_barrier_no_id(__global int * foo,int size)7 void __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
success_barrier_global_id(__global int * foo,int size)21 void __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
success_barrier_local_id(__global int * foo,int size)26 void __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
success_barrier_both_ids(__global int * foo,int size)31 void __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
success_nokernel_barrier_no_id(__global int * foo,int size)37 void 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
success_nokernel_barrier_global_id(__global int * foo,int size)49 void 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
success_nokernel_barrier_local_id(__global int * foo,int size)54 void 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
success_nokernel_barrier_both_ids(__global int * foo,int size)59 void 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
error_barrier_no_id(__global int * foo,int size)67 void __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
success_barrier_global_id(__global int * foo,int size)81 void __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
success_barrier_local_id(__global int * foo,int size)86 void __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
success_barrier_both_ids(__global int * foo,int size)91 void __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
success_nokernel_barrier_no_id(__global int * foo,int size)97 void 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
success_nokernel_barrier_global_id(__global int * foo,int size)109 void 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
success_nokernel_barrier_local_id(__global int * foo,int size)114 void 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
success_nokernel_barrier_both_ids(__global int * foo,int size)119 void 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
error_barrier_no_id(__global int * foo,int size)127 void __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)))
error_barrier_no_id_work_group_size(__global int * foo,int size)142 void __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)))
success_barrier_no_id_work_group_size(__global int * foo,int size)157 void __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
success_barrier_global_id(__global int * foo,int size)169 void __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
success_barrier_local_id(__global int * foo,int size)174 void __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
success_barrier_both_ids(__global int * foo,int size)179 void __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
success_nokernel_barrier_no_id(__global int * foo,int size)185 void 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
success_nokernel_barrier_global_id(__global int * foo,int size)197 void 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
success_nokernel_barrier_local_id(__global int * foo,int size)202 void 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
success_nokernel_barrier_both_ids(__global int * foo,int size)207 void 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
error_barrier_no_id(__global int * foo,int size)215 void __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)))
error_barrier_no_id_work_group_size(__global int * foo,int size)230 void __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)))
success_barrier_no_id_work_group_size(__global int * foo,int size)245 void __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
success_barrier_global_id(__global int * foo,int size)257 void __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
success_barrier_local_id(__global int * foo,int size)262 void __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
success_barrier_both_ids(__global int * foo,int size)267 void __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
success_nokernel_barrier_no_id(__global int * foo,int size)273 void 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
success_nokernel_barrier_global_id(__global int * foo,int size)285 void 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
success_nokernel_barrier_local_id(__global int * foo,int size)290 void 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
success_nokernel_barrier_both_ids(__global int * foo,int size)295 void 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