1 //
2 // Copyright (c) 2017 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 // http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16 #include <stdio.h>
17 #include <string.h>
18 #include "harness/testHarness.h"
19 #include "harness/typeWrappers.h"
20
21 #include <vector>
22
23 #include "procs.h"
24 #include "utils.h"
25 #include <time.h>
26
27
28 #ifdef CL_VERSION_2_0
29
30 static const char* block_global_scope[] =
31 {
32 NL, "int __constant globalVar = 7;"
33 NL, "int (^__constant globalBlock)(int) = ^int(int num)"
34 NL, "{"
35 NL, " return globalVar * num * (1+ get_global_id(0));"
36 NL, "};"
37 NL, "kernel void block_global_scope(__global int* res)"
38 NL, "{"
39 NL, " size_t tid = get_global_id(0);"
40 NL, " res[tid] = -1;"
41 NL, " res[tid] = globalBlock(3) - 21*(tid + 1);"
42 NL, "}"
43 NL
44 };
45
46 static const char* block_kernel_scope[] =
47 {
48 NL, "kernel void block_kernel_scope(__global int* res)"
49 NL, "{"
50 NL, " int multiplier = 3;"
51 NL, " int (^kernelBlock)(int) = ^(int num)"
52 NL, " {"
53 NL, " return num * multiplier;"
54 NL, " };"
55 NL, " size_t tid = get_global_id(0);"
56 NL, " res[tid] = -1;"
57 NL, " multiplier = 8;"
58 NL, " res[tid] = kernelBlock(7) - 21;"
59 NL, "}"
60 NL
61 };
62
63 static const char* block_statement_scope[] =
64 {
65 NL, "kernel void block_statement_scope(__global int* res)"
66 NL, "{"
67 NL, " int multiplier = 0;"
68 NL, " size_t tid = get_global_id(0);"
69 NL, " res[tid] = -1;"
70 NL, " multiplier = 9;"
71 NL, " res[tid] = ^int(int num) { return multiplier * num; } (11) - 99;"
72 NL, "}"
73 NL
74 };
75
76 static const char* block_function_scope[] =
77 {
78 NL, "int fnTest(int a)"
79 NL, "{"
80 NL, " int localVar = 17;"
81 NL, " int (^functionBlock)(int) = ^(int num)"
82 NL, " {"
83 NL, " return localVar * num;"
84 NL, " };"
85 NL, " return 111 - functionBlock(a+1);"
86 NL, "}"
87 NL, "kernel void block_function_scope(__global int* res)"
88 NL, "{"
89 NL, " size_t tid = get_global_id(0);"
90 NL, " res[tid] = -1;"
91 NL, " res[tid] = fnTest(5) - 9;"
92 NL, "}"
93 NL
94 };
95
96 static const char* block_nested_scope[] =
97 {
98 NL, "kernel void block_nested_scope(__global int* res)"
99 NL, "{"
100 NL, " int multiplier = 3;"
101 NL, " int (^kernelBlock)(int) = ^(int num)"
102 NL, " {"
103 NL, " int (^innerBlock)(int) = ^(int n)"
104 NL, " {"
105 NL, " return multiplier * n;"
106 NL, " };"
107 NL, " return num * innerBlock(23);"
108 NL, " };"
109 NL, " size_t tid = get_global_id(0);"
110 NL, " res[tid] = -1;"
111 NL, " multiplier = 8;"
112 NL, " res[tid] = kernelBlock(13) - 897;"
113 NL, "}"
114 NL
115 };
116
117 static const char* block_arg_struct[] =
118 {
119 NL, "struct two_ints {"
120 NL, " short x;"
121 NL, " long y;"
122 NL, "};"
123 NL, "struct two_structs {"
124 NL, " struct two_ints a;"
125 NL, " struct two_ints b;"
126 NL, "};"
127 NL, "kernel void block_arg_struct(__global int* res)"
128 NL, "{"
129 NL, " int (^kernelBlock)(struct two_ints, struct two_structs) = ^int(struct two_ints ti, struct two_structs ts)"
130 NL, " {"
131 NL, " return ti.x * ti.y * ts.a.x * ts.a.y * ts.b.x * ts.b.y;"
132 NL, " };"
133 NL, " struct two_ints i;"
134 NL, " i.x = 2;"
135 NL, " i.y = 3;"
136 NL, " struct two_structs s;"
137 NL, " s.a.x = 4;"
138 NL, " s.a.y = 5;"
139 NL, " s.b.x = 6;"
140 NL, " s.b.y = 7;"
141 NL, " size_t tid = get_global_id(0);"
142 NL, " res[tid] = -1;"
143 NL, " res[tid] = kernelBlock(i,s) - 5040;"
144 NL, "}"
145 NL
146 };
147
148 static const char* block_arg_types_mix[] =
149 {
150 NL, "union number {"
151 NL, " long l;"
152 NL, " float f;"
153 NL, "};"
154 NL, "enum color {"
155 NL, " RED = 0,"
156 NL, " GREEN,"
157 NL, " BLUE" // Using this value - it is actualy "2"
158 NL, "};"
159 NL, "typedef int _INT ;"
160 NL, "typedef char _ACHAR[3] ;"
161 NL, "kernel void block_arg_types_mix(__global int* res)"
162 NL, "{"
163 NL, " int (^kernelBlock)(_INT, _ACHAR, union number, enum color, int, int, int, int, int, int, int, int, int, int, int, int, int) ="
164 NL, " ^int(_INT bi, _ACHAR bch, union number bn, enum color bc, int i1, int i2, int i3, int i4, int i5, int i6, int i7, int i8,"
165 NL, " int i9, int i10, int i11, int i12, int i13)"
166 NL, " {"
167 NL, " return bi * bch[0] * bch[1] * bch[2] * bn.l * bc - i1 - i2 - i3 - i4 - i5 - i6 - i7 - i8 - i9 - i10 - i11 - i12 - i13;"
168 NL, " };"
169 NL, " size_t tid = get_global_id(0);"
170 NL, " res[tid] = -1;"
171 NL, " _INT x = -5;"
172 NL, " _ACHAR char_arr = { 1, 2, 3 };"
173 NL, " union number n;"
174 NL, " n.l = 4;"
175 NL, " enum color c = BLUE;"
176 NL, " res[tid] = kernelBlock(x,char_arr,n,c,1,2,3,4,5,6,7,8,9,10,11,12,13) + 331;"
177 NL, "}"
178 NL
179 };
180
181 static const char* block_arg_pointer[] =
182 {
183 NL, "struct two_ints {"
184 NL, " short x;"
185 NL, " long y;"
186 NL, "};"
187 NL, "kernel void block_arg_pointer(__global int* res)"
188 NL, "{"
189 NL, " int (^kernelBlock)(struct two_ints*, struct two_ints*, int*, int*) = "
190 NL, " ^int(struct two_ints* bs1, struct two_ints* bs2, int* bi1, int* bi2)"
191 NL, " {"
192 NL, " return (*bs1).x * (*bs1).y * (*bs2).x * (*bs2).y * (*bi1) * (*bi2);"
193 NL, " };"
194 NL, " size_t tid = get_global_id(0);"
195 NL, " res[tid] = -1;"
196 NL, " struct two_ints s[2];"
197 NL, " s[0].x = 4;"
198 NL, " s[0].y = 5;"
199 NL, " struct two_ints* ps = s + 1;"
200 NL, " (*ps).x = 6;"
201 NL, " (*ps).y = 7;"
202 NL, " int i = 2;"
203 NL, " int * pi = &i;"
204 NL, " res[tid] = kernelBlock(s,ps,&i,pi) - 3360;"
205 NL, "}"
206 NL
207 };
208
209 static const char* block_arg_global_p[] =
210 {
211 NL, "kernel void block_arg_global_p(__global int* res)"
212 NL, "{"
213 NL, " size_t tid = get_global_id(0);"
214 NL, " res[tid] = -1;"
215 NL, " typedef __global int* int_ptr_to_global_t;"
216 NL, " int_ptr_to_global_t (^kernelBlock)(__global int*, int) =^ int_ptr_to_global_t (__global int* bres, int btid)"
217 NL, " {"
218 NL, " bres[tid] = 5;"
219 NL, " return bres;"
220 NL, " };"
221 NL, " res = kernelBlock(res, tid);"
222 NL, " res[tid] -= 5;"
223 NL, "}"
224 NL
225 };
226
227 static const char* block_arg_const_p[] =
228 {
229 NL, "constant int ci = 8;"
230 NL, "kernel void block_arg_const_p(__global int* res)"
231 NL, "{"
232 NL, " __constant int* (^kernelBlock)(__constant int*) = ^(__constant int* bpci)"
233 NL, " {"
234 NL, " return bpci;"
235 NL, " };"
236 NL, " constant int* pci = &ci;"
237 NL, " constant int* pci_check;"
238 NL, " pci_check = kernelBlock(pci);"
239 NL, " size_t tid = get_global_id(0);"
240 NL, " res[tid] = pci == pci_check ? 0 : -1;"
241 NL, "}"
242 NL
243 };
244
245 static const char* block_ret_struct[] =
246 {
247 NL, "kernel void block_ret_struct(__global int* res)"
248 NL, "{"
249 NL, " struct A {"
250 NL, " int a;"
251 NL, " }; "
252 NL, " struct A (^kernelBlock)(struct A) = ^struct A(struct A a)"
253 NL, " { "
254 NL, " a.a = 6;"
255 NL, " return a;"
256 NL, " };"
257 NL, " size_t tid = get_global_id(0);"
258 NL, " res[tid] = -1;"
259 NL, " struct A aa;"
260 NL, " aa.a = 5;"
261 NL, " res[tid] = kernelBlock(aa).a - 6;"
262 NL, "}"
263 NL
264 };
265
266 static const char* block_arg_global_var[] =
267 {
268 NL, "constant int gi = 8;"
269 NL, "kernel void block_arg_global_var(__global int* res)"
270 NL, "{"
271 NL, " int (^kernelBlock)(int) = ^(int bgi)"
272 NL, " {"
273 NL, " return bgi - 8;"
274 NL, " };"
275 NL, " size_t tid = get_global_id(0);"
276 NL, " res[tid] = kernelBlock(gi);"
277 NL, "}"
278 NL
279 };
280
281 static const char* block_in_for_init[] =
282 {
283 NL, "kernel void block_in_for_init(__global int* res)"
284 NL, "{"
285 NL, " int multiplier = 3;"
286 NL, " int (^kernelBlock)(int) = ^(int num)"
287 NL, " {"
288 NL, " return num * multiplier;"
289 NL, " };"
290 NL, " size_t tid = get_global_id(0);"
291 NL, " res[tid] = 27;"
292 NL, " for(int i=kernelBlock(9); i>0; i--)"
293 NL, " {"
294 NL, " res[tid]--;"
295 NL, " }"
296 NL, "}"
297 NL
298 };
299
300 static const char* block_in_for_cond[] =
301 {
302 NL, "kernel void block_in_for_cond(__global int* res)"
303 NL, "{"
304 NL, " int multiplier = 3;"
305 NL, " int (^kernelBlock)(int) = ^(int num)"
306 NL, " {"
307 NL, " return num * multiplier;"
308 NL, " };"
309 NL, " size_t tid = get_global_id(0);"
310 NL, " res[tid] = 39;"
311 NL, " for(int i=0; i<kernelBlock(13); i++)"
312 NL, " {"
313 NL, " res[tid]--;"
314 NL, " }"
315 NL, "}"
316 NL
317 };
318
319 static const char* block_in_for_iter[] =
320 {
321 NL, "kernel void block_in_for_iter(__global int* res)"
322 NL, "{"
323 NL, " int multiplier = 2;"
324 NL, " int (^kernelBlock)(int) = ^(int num)"
325 NL, " {"
326 NL, " return num * multiplier;"
327 NL, " };"
328 NL, " size_t tid = get_global_id(0);"
329 NL, " res[tid] = 4;"
330 NL, " for(int i=2; i<17; i=kernelBlock(i))"
331 NL, " {"
332 NL, " res[tid]--;"
333 NL, " }"
334 NL, "}"
335 NL
336 };
337
338 static const char* block_in_while_cond[] =
339 {
340 NL, "kernel void block_in_while_cond(__global int* res)"
341 NL, "{"
342 NL, " int (^kernelBlock)(int) = ^(int num)"
343 NL, " {"
344 NL, " return res[num];"
345 NL, " };"
346 NL, " size_t tid = get_global_id(0);"
347 NL, " res[tid] = 27*(tid+1);"
348 NL, " while(kernelBlock(tid))"
349 NL, " {"
350 NL, " res[tid]--;"
351 NL, " }"
352 NL, "}"
353 NL
354 };
355
356 static const char* block_in_while_body[] =
357 {
358 NL, "kernel void block_in_while_body(__global int* res)"
359 NL, "{"
360 NL, " int multiplier = 3;"
361 NL, " int (^kernelBlock)(int) = ^(int num)"
362 NL, " {"
363 NL, " return num * multiplier;"
364 NL, " };"
365 NL, " size_t tid = get_global_id(0);"
366 NL, " int i = 7;"
367 NL, " res[tid] = 3*(7+6+5+4+3+2+1);"
368 NL, " while(i)"
369 NL, " {"
370 NL, " res[tid]-=kernelBlock(i--);"
371 NL, " }"
372 NL, "}"
373 NL
374 };
375
376 static const char* block_in_do_while_body[] =
377 {
378 NL, "kernel void block_in_do_while_body(__global int* res)"
379 NL, "{"
380 NL, " int multiplier = 3;"
381 NL, " size_t tid = get_global_id(0);"
382 NL, " int i = 100;"
383 NL, " res[tid] = 3*5050;"
384 NL, " do"
385 NL, " {"
386 NL, " int (^kernelBlock)(int) = ^(int num)"
387 NL, " {"
388 NL, " return num * multiplier;"
389 NL, " };"
390 NL, " res[tid]-=kernelBlock(i--);"
391 NL, " } while(i);"
392 NL, "}"
393 NL
394 };
395
396 static const char* block_cond_statement[] =
397 {
398 NL, "kernel void block_cond_statement(__global int* res)"
399 NL, "{"
400 NL, " int multiplier = 2;"
401 NL, " int (^kernelBlock)(int) = ^(int num)"
402 NL, " {"
403 NL, " return num * multiplier;"
404 NL, " };"
405 NL, " size_t tid = get_global_id(0);"
406 NL, " res[tid] = 120;"
407 NL, " res[tid] = (kernelBlock(2) == 4) ? res[tid] - 30 : res[tid] - 1;"
408 NL, " res[tid] = (kernelBlock(2) == 5) ? res[tid] - 3 : res[tid] - 30;"
409 NL, " res[tid] = (1) ? res[tid] - kernelBlock(15) : res[tid] - 7;"
410 NL, " res[tid] = (0) ? res[tid] - 13 : res[tid] - kernelBlock(15);"
411 NL, "}"
412 NL
413 };
414
415 static const char* block_in_if_cond[] =
416 {
417 NL, "kernel void block_in_if_cond(__global int* res)"
418 NL, "{"
419 NL, " int multiplier = 2;"
420 NL, " int (^kernelBlock)(int) = ^(int num)"
421 NL, " {"
422 NL, " return num * multiplier;"
423 NL, " };"
424 NL, " size_t tid = get_global_id(0);"
425 NL, " res[tid] = 7;"
426 NL, " if (kernelBlock(5))"
427 NL, " {"
428 NL, " res[tid]-= 3;"
429 NL, " }"
430 NL, " if (kernelBlock(0))"
431 NL, " {"
432 NL, " res[tid]-= 2;"
433 NL, " }"
434 NL, " else"
435 NL, " {"
436 NL, " res[tid]-= 4;"
437 NL, " }"
438 NL, "}"
439 NL
440 };
441
442 static const char* block_in_if_branch[] =
443 {
444 NL, "kernel void block_in_if_branch(__global int* res)"
445 NL, "{"
446 NL, " int multiplier = 2;"
447 NL, " int (^kernelBlock)(int) = ^(int num)"
448 NL, " {"
449 NL, " return num * multiplier;"
450 NL, " };"
451 NL, " size_t tid = get_global_id(0);"
452 NL, " res[tid] = 7;"
453 NL, " if (kernelBlock(5))"
454 NL, " {"
455 NL, " res[tid]-= ^(int num){ return num - 1; }(4);" // res[tid]-=3;
456 NL, " }"
457 NL, " if (kernelBlock(0))"
458 NL, " {"
459 NL, " res[tid]-= ^(int num){ return num - 1; }(3);" // res[tid]-=2;
460 NL, " }"
461 NL, " else"
462 NL, " {"
463 NL, " int (^ifBlock)(int) = ^(int num)"
464 NL, " {"
465 NL, " return num + 1;"
466 NL, " };"
467 NL, " res[tid]-= ifBlock(3);" // res[tid]-=4;
468 NL, " }"
469 NL, "}"
470 NL
471 };
472
473 static const char* block_switch_cond[] =
474 {
475 NL, "kernel void block_switch_cond(__global int* res)"
476 NL, "{"
477 NL, " int multiplier = 2;"
478 NL, " int (^kernelBlock)(int) = ^(int num)"
479 NL, " {"
480 NL, " return num * multiplier;"
481 NL, " };"
482 NL, " size_t tid = get_global_id(0);"
483 NL, " res[tid] = 12;"
484 NL, " int i = 1;"
485 NL, " while(i <= 3)"
486 NL, " {"
487 NL, " switch (kernelBlock(i))"
488 NL, " {"
489 NL, " case 2:"
490 NL, " res[tid] = res[tid] - 2;"
491 NL, " break;"
492 NL, " case 4:"
493 NL, " res[tid] = res[tid] - 4;"
494 NL, " break;"
495 NL, " case 6:"
496 NL, " res[tid] = res[tid] - 6;"
497 NL, " break;"
498 NL, " default:"
499 NL, " break;"
500 NL, " }"
501 NL, " i++;"
502 NL, " }"
503 NL, "}"
504 NL
505 };
506
507 static const char* block_switch_case[] =
508 {
509 NL, "kernel void block_switch_case(__global int* res)"
510 NL, "{"
511 NL, " int multiplier = 2;"
512 NL, " int (^kernelBlock)(int) = ^(int num)"
513 NL, " {"
514 NL, " return num * multiplier;"
515 NL, " };"
516 NL, " size_t tid = get_global_id(0);"
517 NL, " res[tid] = 12;"
518 NL, " int i = 1;"
519 NL, " while(i <= 3)"
520 NL, " {"
521 NL, " switch (kernelBlock(i))"
522 NL, " {"
523 NL, " case 2:"
524 NL, " res[tid]-=^(int num){ return num - 1; }(3);" // res[tid]-=2;
525 NL, " break;"
526 NL, " case 4:"
527 NL, " {"
528 NL, " int (^caseBlock)(int) = ^(int num)"
529 NL, " {"
530 NL, " return num + 1;"
531 NL, " };"
532 NL, " res[tid]-=caseBlock(3);" // res[tid]-=4;
533 NL, " break;"
534 NL, " }"
535 NL, " case 6:"
536 NL, " res[tid]-=kernelBlock(3);" // res[tid]-=6;
537 NL, " break;"
538 NL, " default:"
539 NL, " break;"
540 NL, " }"
541 NL, " i++;"
542 NL, " }"
543 NL, "}"
544 NL
545 };
546
547 // Accessing data from Block
548
549 static const char* block_access_program_data[] =
550 {
551 NL, "int __constant globalVar1 = 7;"
552 NL, "int __constant globalVar2 = 11;"
553 NL, "int __constant globalVar3 = 13;"
554 NL, "int (^__constant globalBlock)(int) = ^int(int num)"
555 NL, "{"
556 NL, " return globalVar1 * num;"
557 NL, "};"
558 NL, "kernel void block_access_program_data(__global int* res)"
559 NL, "{"
560 NL, " int (^ kernelBlock)(int) = ^int(int num)"
561 NL, " {"
562 NL, " return globalVar2 * num;"
563 NL, " };"
564 NL, " size_t tid = get_global_id(0);"
565 NL, " res[tid] = tid + 1;"
566 NL, " res[tid] = globalBlock(res[tid]);"
567 NL, " res[tid] = kernelBlock(res[tid]);"
568 NL, " res[tid] = ^(int num){ return globalVar3*num; }(res[tid]) - (7*11*13)*(tid + 1);"
569 NL, "}"
570 NL
571 };
572
573 static const char* block_access_kernel_data[] =
574 {
575 NL, "kernel void block_access_kernel_data(__global int* res)"
576 NL, "{"
577 NL, " int var1 = 7;"
578 NL, " int var2 = 11;"
579 NL, " int var3 = 13;"
580 NL, " int (^ kernelBlock)(int) = ^int(int num)"
581 NL, " {"
582 NL, " int (^ nestedBlock)(int) = ^int (int num)"
583 NL, " {"
584 NL, " return var1 * num;"
585 NL, " };"
586 NL, " return var2 * nestedBlock(num);"
587 NL, " };"
588 NL, " size_t tid = get_global_id(0);"
589 NL, " res[tid] = tid + 1;"
590 NL, " res[tid] = kernelBlock(res[tid]);"
591 NL, " res[tid] = ^(int num){ return var3*num; }(res[tid]) - (7*11*13)*(tid + 1);"
592 NL, "}"
593 NL
594 };
595
596 static const char* block_access_chained_data[] =
597 {
598 NL, "kernel void block_access_chained_data(__global int* res)"
599 NL, "{"
600 NL, " int (^ kernelBlock)(int) = ^int(int num)"
601 NL, " {"
602 NL, " int var1 = 7;"
603 NL, " int var2 = 11;"
604 NL, " int var3 = 13;"
605 NL, " int (^ nestedBlock1)(int) = ^int (int num)"
606 NL, " {"
607 NL, " int (^ nestedBlock2) (int) = ^int (int num)"
608 NL, " {"
609 NL, " return var2 * ^(int num){ return var3*num; }(num);"
610 NL, " };"
611 NL, " return var1 * nestedBlock2(num);"
612 NL, " };"
613 NL, " return nestedBlock1(num);"
614 NL, " };"
615 NL, " size_t tid = get_global_id(0);"
616 NL, " res[tid] = tid + 1;"
617 NL, " res[tid] = kernelBlock(res[tid]) - (7*11*13)*(tid + 1);"
618 NL, "}"
619 NL
620 };
621
622 static const char* block_access_volatile_data[] =
623 {
624 NL, "kernel void block_access_volatile_data(__global int* res)"
625 NL, "{"
626 NL, " int var1 = 7;"
627 NL, " int var2 = 11;"
628 NL, " volatile int var3 = 13;"
629 NL, ""
630 NL, " int (^ kernelBlock)(int) = ^int(int num)"
631 NL, " {"
632 NL, " int (^ nestedBlock)(int) = ^int (int num)"
633 NL, " {"
634 NL, " return var1 * num;"
635 NL, " };"
636 NL, " return var2 * nestedBlock(num);"
637 NL, " };"
638 NL, " size_t tid = get_global_id(0);"
639 NL, " res[tid] = tid + 1;"
640 NL, " res[tid] = kernelBlock(res[tid]);"
641 NL, " res[tid] = ^(int num){ return var3*num; }(res[tid]) - (7*11*13)*(tid + 1);"
642 NL, "}"
643 NL
644 };
645
646 static const char* block_typedef_kernel[] =
647 {
648 NL, "kernel void block_typedef_kernel(__global int* res)"
649 NL, "{"
650 NL, " typedef int* (^block_t)(int*);"
651 NL, " size_t tid = get_global_id(0);"
652 NL, " res[tid] = -1;"
653 NL, " int i[4] = { 3, 4, 4, 1 };"
654 NL, " int *temp = i; // workaround clang bug"
655 NL, " block_t kernelBlock = ^(int* pi)"
656 NL, " {"
657 NL, " block_t b = ^(int* n) { return n - 1; };"
658 NL, " return pi + *(b(temp+4));"
659 NL, " };"
660 NL, " switch (*(kernelBlock(i))) {"
661 NL, " case 4:"
662 NL, " res[tid] += *(kernelBlock(i+1));"
663 NL, " break;"
664 NL, " default:"
665 NL, " res[tid] = -100;"
666 NL, " break;"
667 NL, " }"
668 NL, " res[tid] += *(kernelBlock(i)) - 7;"
669 NL, "}"
670 NL
671 };
672
673 static const char* block_typedef_func[] =
674 {
675 NL, "int func(int fi)"
676 NL, "{"
677 NL, " typedef int (^block_t)(int);"
678 NL, " const block_t funcBlock = ^(int bi)"
679 NL, " {"
680 NL, " typedef short (^block2_t)(short);"
681 NL, " block2_t nestedBlock = ^(short ni)"
682 NL, " {"
683 NL, " return (short)(ni - 1);"
684 NL, " };"
685 NL, " return bi * nestedBlock(3);"
686 NL, " };"
687 NL, " return funcBlock(fi * 2);"
688 NL, "}"
689 NL, "kernel void block_typedef_func(__global int* res)"
690 NL, "{"
691 NL, " size_t tid = get_global_id(0);"
692 NL, " res[tid] = -1;"
693 NL, " res[tid] = func(1) - 4;"
694 NL, "}"
695 NL
696 };
697
698 static const char* block_typedef_stmnt_if[] =
699 {
700 NL, "kernel void block_typedef_stmnt_if(__global int* res)"
701 NL, "{ "
702 NL, " int flag = 1;"
703 NL, " int sum = 0;"
704 NL, " if (flag) {"
705 NL, " typedef int (^block_t)(int);"
706 NL, " const block_t kernelBlock = ^(int bi)"
707 NL, " {"
708 NL, " block_t b = ^(int bi)"
709 NL, " {"
710 NL, " return bi + 1;"
711 NL, " };"
712 NL, " return bi + b(1);"
713 NL, " };"
714 NL, " sum = kernelBlock(sum);"
715 NL, " }"
716 NL, " size_t tid = get_global_id(0);"
717 NL, " res[tid] = sum - 2;"
718 NL, "}"
719 NL
720 };
721
722 static const char* block_typedef_loop[] =
723 {
724 NL, "kernel void block_typedef_loop(__global int* res)"
725 NL, "{ "
726 NL, " int sum = -1;"
727 NL, " for (int i = 0; i < 3; i++) {"
728 NL, " typedef int (^block_t)(void);"
729 NL, " const block_t kernelBlock = ^()"
730 NL, " {"
731 NL, " return i + 1;"
732 NL, " };"
733 NL, " sum += kernelBlock();"
734 NL, " }"
735 NL, " size_t tid = get_global_id(0);"
736 NL, " res[tid] = sum - 5;"
737 NL, "}"
738 NL
739 };
740
741 static const char* block_typedef_mltpl_func[] =
742 {
743 NL, "int func(int fi)"
744 NL, "{"
745 NL, " typedef int (^block_t)(int);"
746 NL, " typedef int (^block2_t)(int);"
747 NL, " const block_t funcBlock1 = ^(int bi) { return bi; };"
748 NL, " const block2_t funcBlock2 = ^(int bi)"
749 NL, " {"
750 NL, " typedef short (^block3_t)(short);"
751 NL, " typedef short (^block4_t)(short);"
752 NL, " const block3_t nestedBlock1 = ^(short ni)"
753 NL, " {"
754 NL, " return (short)(ni - 1);"
755 NL, " };"
756 NL, " const block4_t nestedBlock2 = ^(short ni)"
757 NL, " {"
758 NL, " return (short)(ni - 2);"
759 NL, " };"
760 NL, " return bi * nestedBlock1(3) * nestedBlock2(3);"
761 NL, " };"
762 NL, " return funcBlock2(fi * 2) + funcBlock1(1);"
763 NL, "}"
764 NL, "kernel void block_typedef_mltpl_func(__global int* res)"
765 NL, "{"
766 NL, " size_t tid = get_global_id(0);"
767 NL, " res[tid] = -1;"
768 NL, " typedef int (^block1_t)(int);"
769 NL, " typedef int (^block2_t)(int);"
770 NL, " const block1_t kernelBlock1 = ^(int bi) { return bi + 8; };"
771 NL, " const block2_t kernelBlock2 = ^(int bi) { return bi + 3; };"
772 NL, " res[tid] = func(1) - kernelBlock1(2) / kernelBlock2(-1);"
773 NL, "}"
774 NL
775 };
776
777 static const char* block_typedef_mltpl_stmnt[] =
778 {
779 NL, "kernel void block_typedef_mltpl_stmnt(__global int* res)"
780 NL, "{"
781 NL, " size_t tid = get_global_id(0);"
782 NL, " res[tid] = -1;"
783 NL, " int a;"
784 NL, " do"
785 NL, " {"
786 NL, " typedef float (^blockf_t)(float);"
787 NL, " typedef int (^blocki_t)(int);"
788 NL, " const blockf_t blockF = ^(float bi) { return (float)(bi + 3.3); };"
789 NL, " const blocki_t blockI = ^(int bi) { return bi + 2; };"
790 NL, " if ((blockF(.0)-blockI(0)) > 0)"
791 NL, " {"
792 NL, " typedef uint (^block_t)(uint);"
793 NL, " const block_t nestedBlock = ^(uint bi) { return (uint)(bi + 4); };"
794 NL, " a = nestedBlock(1) + nestedBlock(2);"
795 NL, " break;"
796 NL, " }"
797 NL, " } while(1); "
798 NL, " res[tid] = a - 11;"
799 NL, "}"
800 NL
801 };
802
803 static const char* block_typedef_mltpl_g[] =
804 {
805 NL, "typedef int (^block1_t)(float, int); "
806 NL, "constant block1_t b1 = ^(float fi, int ii) { return (int)(ii + fi); };"
807 NL, "typedef int (^block2_t)(float, int);"
808 NL, "constant block2_t b2 = ^(float fi, int ii) { return (int)(ii + fi); };"
809 NL, "typedef float (^block3_t)(int, int);"
810 NL, "constant block3_t b3 = ^(int i1, int i2) { return (float)(i1 + i2); };"
811 NL, "typedef int (^block4_t)(float, float);"
812 NL, "kernel void block_typedef_mltpl_g(__global int* res)"
813 NL, "{"
814 NL, " size_t tid = get_global_id(0);"
815 NL, " res[tid] = -1;"
816 NL, " block4_t b4 = ^(float f1, float f2) { return (int)(f1 + f2); };"
817 NL, " res[tid] = b1(1.1, b2(1.1, 1)) - b4(b3(1,1), 1.1);"
818 NL, "}"
819 NL
820 };
821
822 static const char* block_literal[] =
823 {
824 NL, "int func()"
825 NL, "{"
826 NL, " return ^(int i) {"
827 NL, " return ^(ushort us)"
828 NL, " {"
829 NL, " return (int)us + i;"
830 NL, " }(3);"
831 NL, " }(7) - 10;"
832 NL, "}"
833 NL, "kernel void block_literal(__global int* res)"
834 NL, "{"
835 NL, " size_t tid = get_global_id(0);"
836 NL, " res[tid] = -1;"
837 NL, " res[tid] = func();"
838 NL, "}"
839 NL
840 };
841
842 static const char* block_complex[] =
843 {
844 NL, "kernel void block_complex(__global int* res)"
845 NL, "{"
846 NL, " int (^kernelBlock)(int) = ^(int num)"
847 NL, " {"
848 NL, " int result = 1;"
849 NL, " for (int i = 0; i < num; i++)"
850 NL, " {"
851 NL, " switch(i)"
852 NL, " {"
853 NL, " case 0:"
854 NL, " case 1:"
855 NL, " case 2:"
856 NL, " result += i;"
857 NL, " break;"
858 NL, " case 3:"
859 NL, " if (result < num)"
860 NL, " result += i;"
861 NL, " else"
862 NL, " result += i * 2;"
863 NL, " break;"
864 NL, " case 4:"
865 NL, " while (1)"
866 NL, " {"
867 NL, " result++;"
868 NL, " if (result)"
869 NL, " goto ret;"
870 NL, " }"
871 NL, " break;"
872 NL, " default:"
873 NL, " return 777;"
874 NL, " }"
875 NL, " }"
876 NL, " ret: ;"
877 NL, " while (num) {"
878 NL, " num--;"
879 NL, " if (num % 2 == 0)"
880 NL, " continue;"
881 NL, " result++;"
882 NL, " }"
883 NL, " return result;"
884 NL, " };"
885 NL, " size_t tid = get_global_id(0);"
886 NL, " res[tid] = -1;"
887 NL, " res[tid] = kernelBlock(7) - 11;"
888 NL, "}"
889 NL
890 };
891
892 static const char* block_empty[] =
893 {
894 NL, "kernel void block_empty(__global int* res)"
895 NL, "{"
896 NL, " void (^kernelBlock)(void) = ^(){};"
897 NL, " size_t tid = get_global_id(0);"
898 NL, " res[tid] = -1;"
899 NL, " kernelBlock();"
900 NL, " res[tid] = 0;"
901 NL, "}"
902 NL
903 };
904
905 static const char* block_builtin[] =
906 {
907 NL, "kernel void block_builtin(__global int* res)"
908 NL, "{"
909 NL, " int b = 3;"
910 NL, " int (^kernelBlock)(int) = ^(int a)"
911 NL, " {"
912 NL, " return (int)abs(a - b);"
913 NL, " };"
914 NL, " size_t tid = get_global_id(0);"
915 NL, " res[tid] = -1;"
916 NL, " res[tid] = kernelBlock(2) - 1;"
917 NL, "}"
918 NL
919 };
920
921 static const char* block_barrier[] =
922 {
923 NL, "kernel void block_barrier(__global int* res)"
924 NL, "{"
925 NL, " int b = 3;"
926 NL, " size_t tid = get_global_id(0);"
927 NL, " size_t lsz = get_local_size(0);"
928 NL, " size_t gid = get_group_id(0);"
929 NL, " size_t idx = gid*lsz;"
930 NL, ""
931 NL, " res[tid]=lsz;"
932 NL, " barrier(CLK_GLOBAL_MEM_FENCE);"
933 NL, " int (^kernelBlock)(int) = ^(int a)"
934 NL, " {"
935 NL, " atomic_dec(res+idx);"
936 NL, " barrier(CLK_GLOBAL_MEM_FENCE);"
937 NL, " return (int)abs(a - b) - (res[idx] != 0 ? 0 : 1);"
938 NL, " };"
939 NL, ""
940 NL, " int d = kernelBlock(2);"
941 NL, " res[tid] = d;"
942 NL, "}"
943 NL
944 };
945
946
947
948 static const kernel_src sources_execute_block[] =
949 {
950 // Simple blocks
951 KERNEL(block_global_scope),
952 KERNEL(block_kernel_scope),
953 KERNEL(block_statement_scope),
954 KERNEL(block_function_scope),
955 KERNEL(block_nested_scope),
956
957 // Kernels with Block in for/while/if/switch
958 KERNEL(block_in_for_init),
959 KERNEL(block_in_for_cond),
960 KERNEL(block_in_for_iter),
961 KERNEL(block_in_while_cond),
962 KERNEL(block_in_while_body),
963 KERNEL(block_in_do_while_body),
964 KERNEL(block_cond_statement),
965 KERNEL(block_in_if_cond),
966 KERNEL(block_in_if_branch),
967 KERNEL(block_switch_cond),
968 KERNEL(block_switch_case),
969 KERNEL(block_literal),
970
971 // Accessing data from block
972 KERNEL(block_access_program_data),
973 KERNEL(block_access_kernel_data),
974 KERNEL(block_access_chained_data),
975 KERNEL(block_access_volatile_data),
976
977 // Block args
978 KERNEL(block_arg_struct),
979 KERNEL(block_arg_types_mix),
980 KERNEL(block_arg_pointer),
981 KERNEL(block_arg_global_p),
982 KERNEL(block_arg_const_p),
983 KERNEL(block_ret_struct),
984 KERNEL(block_arg_global_var),
985
986 // Block in typedef
987 KERNEL(block_typedef_kernel),
988 KERNEL(block_typedef_func),
989 KERNEL(block_typedef_stmnt_if),
990 KERNEL(block_typedef_loop),
991 KERNEL(block_typedef_mltpl_func),
992 KERNEL(block_typedef_mltpl_stmnt),
993 KERNEL(block_typedef_mltpl_g),
994
995 // Non - trivial blocks
996 KERNEL(block_complex),
997 KERNEL(block_empty),
998 KERNEL(block_builtin),
999 KERNEL(block_barrier),
1000
1001 };
1002 static const size_t num_kernels_execute_block = arr_size(sources_execute_block);
1003
check_kernel_results(cl_int * results,cl_int len)1004 static int check_kernel_results(cl_int* results, cl_int len)
1005 {
1006 for(cl_int i = 0; i < len; ++i)
1007 {
1008 if(results[i] != 0) return i;
1009 }
1010 return -1;
1011 }
1012
test_execute_block(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)1013 int test_execute_block(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
1014 {
1015 size_t i;
1016 size_t ret_len;
1017 cl_int n, err_ret, res = 0;
1018 clCommandQueueWrapper dev_queue;
1019 cl_int kernel_results[MAX_GWS] = {0xDEADBEEF};
1020
1021 size_t max_local_size = 1;
1022 err_ret = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(max_local_size), &max_local_size, &ret_len);
1023 test_error(err_ret, "clGetDeviceInfo(CL_DEVICE_MAX_WORK_GROUP_SIZE) failed");
1024
1025 size_t global_size = MAX_GWS;
1026 size_t local_size = (max_local_size > global_size/16) ? global_size/16 : max_local_size;
1027
1028 size_t failCnt = 0;
1029 for(i = 0; i < num_kernels_execute_block; ++i)
1030 {
1031 if (!gKernelName.empty() && gKernelName != sources_execute_block[i].kernel_name)
1032 continue;
1033
1034 log_info("Running '%s' kernel (%d of %d) ...\n", sources_execute_block[i].kernel_name, i + 1, num_kernels_execute_block);
1035 err_ret = run_n_kernel_args(context, queue, sources_execute_block[i].lines, sources_execute_block[i].num_lines, sources_execute_block[i].kernel_name, local_size, global_size, kernel_results, sizeof(kernel_results), 0, NULL);
1036 if(check_error(err_ret, "'%s' kernel execution failed", sources_execute_block[i].kernel_name)) { ++failCnt; res = -1; }
1037 else if((n = check_kernel_results(kernel_results, arr_size(kernel_results))) >= 0 && check_error(-1, "'%s' kernel results validation failed: [%d] returned %d expected 0", sources_execute_block[i].kernel_name, n, kernel_results[n])) { ++failCnt; res = -1; }
1038 else log_info("'%s' kernel is OK.\n", sources_execute_block[i].kernel_name);
1039 }
1040
1041 if (failCnt > 0)
1042 {
1043 log_error("ERROR: %d of %d kernels failed.\n", failCnt, num_kernels_execute_block);
1044 }
1045
1046 return res;
1047 }
1048
1049
1050 #endif
1051
1052