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 "harness/testHarness.h"
17 #include "harness/typeWrappers.h"
18 #include "base.h"
19
20 #include <string>
21 #include <vector>
22 #include <algorithm>
23
24 class CBasicTest : CTest {
25 public:
CBasicTest(const std::vector<std::string> & kernel)26 CBasicTest(const std::vector<std::string>& kernel) : CTest(), _kernels(kernel) {
27
28 }
29
CBasicTest(const std::string & kernel)30 CBasicTest(const std::string& kernel) : CTest(), _kernels(1, kernel) {
31
32 }
33
ExecuteSubcase(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements,const std::string & src)34 int ExecuteSubcase(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, const std::string& src) {
35 cl_int error;
36
37 clProgramWrapper program;
38 clKernelWrapper kernel;
39
40 const char *srcPtr = src.c_str();
41
42 if (create_single_kernel_helper(context, &program, &kernel, 1, &srcPtr,
43 "testKernel"))
44 {
45 log_error("create_single_kernel_helper failed");
46 return -1;
47 }
48
49 size_t bufferSize = num_elements * sizeof(cl_uint);
50 clMemWrapper buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, bufferSize, NULL, &error);
51 test_error(error, "clCreateBuffer failed");
52
53 error = clSetKernelArg(kernel, 0, sizeof(buffer), &buffer);
54 test_error(error, "clSetKernelArg failed");
55
56 size_t globalWorkGroupSize = num_elements;
57 size_t localWorkGroupSize = 0;
58 error = get_max_common_work_group_size(context, kernel, globalWorkGroupSize, &localWorkGroupSize);
59 test_error(error, "Unable to get common work group size");
60
61 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalWorkGroupSize, &localWorkGroupSize, 0, NULL, NULL);
62 test_error(error, "clEnqueueNDRangeKernel failed");
63
64 // verify results
65 std::vector<cl_uint> results(num_elements);
66
67 error = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, bufferSize, &results[0], 0, NULL, NULL);
68 test_error(error, "clEnqueueReadBuffer failed");
69
70 size_t passCount = std::count(results.begin(), results.end(), 1);
71 if (passCount != results.size()) {
72 std::vector<cl_uint>::iterator iter = std::find(results.begin(), results.end(), 0);
73 log_error("Verification on device failed at index %ld\n", std::distance(results.begin(), iter));
74 log_error("%ld out of %ld failed\n", (results.size()-passCount), results.size());
75 return -1;
76 }
77
78 return CL_SUCCESS;
79 }
80
Execute(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)81 int Execute(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
82 cl_int result = CL_SUCCESS;
83
84 for (std::vector<std::string>::const_iterator it = _kernels.begin(); it != _kernels.end(); ++it) {
85 log_info("Executing subcase #%ld out of %ld\n", (it - _kernels.begin() + 1), _kernels.size());
86
87 result |= ExecuteSubcase(deviceID, context, queue, num_elements, *it);
88 }
89
90 return result;
91 }
92
93 private:
94 const std::vector<std::string> _kernels;
95 };
96
test_function_get_fence(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)97 int test_function_get_fence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
98 const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
99 NL
100 NL "__global int gint = 1;"
101 NL "__global uchar guchar = 3;"
102 NL
103 NL "bool helperFunction(int *intp, float *floatp, uchar *ucharp, ushort *ushortp, long *longp) {"
104 NL " if (!isFenceValid(get_fence(intp)))"
105 NL " return false;"
106 NL " if (!isFenceValid(get_fence(floatp)))"
107 NL " return false;"
108 NL " if (!isFenceValid(get_fence(ucharp)))"
109 NL " return false;"
110 NL " if (!isFenceValid(get_fence(ushortp)))"
111 NL " return false;"
112 NL " if (!isFenceValid(get_fence(longp)))"
113 NL " return false;"
114 NL
115 NL " if (*intp != 1 || *floatp != 2.0f || *ucharp != 3 || *ushortp != 4 || *longp != 5)"
116 NL " return false;"
117 NL
118 NL " return true;"
119 NL "}"
120 NL
121 NL "__kernel void testKernel(__global uint *results) {"
122 NL " uint tid = get_global_id(0);"
123 NL
124 NL " __local float lfloat;"
125 NL " lfloat = 2.0f;"
126 NL " __local ushort lushort;"
127 NL " lushort = 4;"
128 NL " long plong = 5;"
129 NL
130 NL " __global int *gintp = &gint;"
131 NL " __local float *lfloatp = &lfloat;"
132 NL " __global uchar *gucharp = &guchar;"
133 NL " __local ushort *lushortp = &lushort;"
134 NL " __private long *plongp = &plong;"
135 NL
136 NL " results[tid] = helperFunction(gintp, lfloatp, gucharp, lushortp, plongp);"
137 NL "}"
138 NL;
139
140 CBasicTest test(KERNEL_FUNCTION);
141
142 return test.Execute(deviceID, context, queue, num_elements);
143 }
144
test_function_to_address_space(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)145 int test_function_to_address_space(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
146 const std::string KERNEL_FUNCTION =
147 NL
148 NL "__global int gint = 1;"
149 NL "__global uchar guchar = 3;"
150 NL
151 NL "bool helperFunction(int *gintp, float *lfloatp, uchar *gucharp, ushort *lushortp, long *plongp) {"
152 NL " if (to_global(gintp) == NULL)"
153 NL " return false;"
154 NL " if (to_local(lfloatp) == NULL)"
155 NL " return false;"
156 NL " if (to_global(gucharp) == NULL)"
157 NL " return false;"
158 NL " if (to_local(lushortp) == NULL)"
159 NL " return false;"
160 NL " if (to_private(plongp) == NULL)"
161 NL " return false;"
162 NL
163 NL " if (*gintp != 1 || *lfloatp != 2.0f || *gucharp != 3 || *lushortp != 4 || *plongp != 5)"
164 NL " return false;"
165 NL
166 NL " return true;"
167 NL "}"
168 NL
169 NL "__kernel void testKernel(__global uint *results) {"
170 NL " uint tid = get_global_id(0);"
171 NL
172 NL " __local float lfloat;"
173 NL " lfloat = 2.0f;"
174 NL " __local ushort lushort;"
175 NL " lushort = 4;"
176 NL " long plong = 5;"
177 NL
178 NL " __global int *gintp = &gint;"
179 NL " __local float *lfloatp = &lfloat;"
180 NL " __global uchar *gucharp = &guchar;"
181 NL " __local ushort *lushortp = &lushort;"
182 NL " __private long *plongp = &plong;"
183 NL
184 NL " results[tid] = helperFunction(gintp, lfloatp, gucharp, lushortp, plongp);"
185 NL "}"
186 NL;
187
188 CBasicTest test(KERNEL_FUNCTION);
189
190 return test.Execute(deviceID, context, queue, num_elements);
191 }
192
test_variable_get_fence(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)193 int test_variable_get_fence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
194 const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
195 NL
196 NL "__global int gint = 1;"
197 NL
198 NL "__kernel void testKernel(__global uint *results) {"
199 NL " uint tid = get_global_id(0);"
200 NL
201 NL " __local ushort lushort;"
202 NL " lushort = 2;"
203 NL " float pfloat = 3.0f;"
204 NL
205 NL " // tested pointers"
206 NL " __global int *gintp = &gint;"
207 NL " __local ushort *lushortp = &lushort;"
208 NL " __private float *pfloatp = &pfloat;"
209 NL
210 NL " int failures = 0;"
211 NL " if (!isFenceValid(get_fence(gintp)))"
212 NL " failures++;"
213 NL " if (!isFenceValid(get_fence(lushortp)))"
214 NL " failures++;"
215 NL " if (!isFenceValid(get_fence(pfloatp)))"
216 NL " failures++;"
217 NL " results[tid] = (failures == 0);"
218 NL "}"
219 NL;
220
221 CBasicTest test(KERNEL_FUNCTION);
222
223 return test.Execute(deviceID, context, queue, num_elements);
224 }
225
test_variable_to_address_space(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)226 int test_variable_to_address_space(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
227 const std::string KERNEL_FUNCTION =
228 NL
229 NL "__global int gint = 1;"
230 NL
231 NL "__kernel void testKernel(__global uint *results) {"
232 NL " uint tid = get_global_id(0);"
233 NL
234 NL " __local ushort lushort;"
235 NL " lushort = 2;"
236 NL " float pfloat = 3.0f;"
237 NL
238 NL " // tested pointers"
239 NL " __global int * gintp = &gint;"
240 NL " __local ushort *lushortp = &lushort;"
241 NL " __private float *pfloatp = &pfloat;"
242 NL
243 NL " int failures = 0;"
244 NL " if (to_global(gintp) == NULL)"
245 NL " failures++;"
246 NL " if (to_local(lushortp) == NULL)"
247 NL " failures++;"
248 NL " if (to_private(pfloatp) == NULL)"
249 NL " failures++;"
250 NL " results[tid] = (failures == 0);"
251 NL "}"
252 NL;
253
254 CBasicTest test(KERNEL_FUNCTION);
255
256 return test.Execute(deviceID, context, queue, num_elements);
257 }
258
test_casting(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)259 int test_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
260 std::vector<std::string> KERNEL_FUNCTIONS;
261
262 // pointers to global, local or private are implicitly convertible to generic
263 KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
264 NL
265 NL "__global int gint = 1;"
266 NL
267 NL "__kernel void testKernel(__global uint *results) {"
268 NL " uint tid = get_global_id(0);"
269 NL
270 NL " __local int lint;"
271 NL " lint = 2;"
272 NL " int pint = 3;"
273 NL
274 NL " // count mismatches with expected fence types"
275 NL " int failures = 0;"
276 NL
277 NL " // tested pointer"
278 NL " // generic can be reassigned to different named address spaces"
279 NL " int * intp;"
280 NL
281 NL " intp = &gint;"
282 NL " failures += !(isFenceValid(get_fence(intp)));"
283 NL " failures += !(to_global(intp));"
284 NL " failures += (*intp != 1);"
285 NL
286 NL " intp = &lint;"
287 NL " failures += !(isFenceValid(get_fence(intp)));"
288 NL " failures += !(to_local(intp));"
289 NL " failures += (*intp != 2);"
290 NL
291 NL " intp = &pint;"
292 NL " failures += !(isFenceValid(get_fence(intp)));"
293 NL " failures += !(to_private(intp));"
294 NL " failures += (*intp != 3);"
295 NL
296 NL " results[tid] = (failures == 0);"
297 NL "}"
298 NL
299 );
300
301 // converting from a generic pointer to a named address space is legal only with explicit casting
302 KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
303 NL
304 NL "__global int gint = 1;"
305 NL
306 NL "__kernel void testKernel(__global uint *results) {"
307 NL " uint tid = get_global_id(0);"
308 NL
309 NL " __local int lint;"
310 NL " lint = 2;"
311 NL " int pint = 3;"
312 NL
313 NL " // count mismatches with expected fence types"
314 NL " int failures = 0;"
315 NL
316 NL " // tested pointer"
317 NL " // generic can be reassigned to different named address spaces"
318 NL " int * intp;"
319 NL
320 NL " intp = &gint;"
321 NL " global int * gintp = (global int *)intp;"
322 NL " failures += !(isFenceValid(get_fence(gintp)));"
323 NL " failures += !(to_global(gintp));"
324 NL " failures += (*gintp != 1);"
325 NL
326 NL " intp = &lint;"
327 NL " local int * lintp = (local int *)intp;"
328 NL " failures += !(isFenceValid(get_fence(lintp)));"
329 NL " failures += !(to_local(lintp));"
330 NL " failures += (*lintp != 2);"
331 NL
332 NL " intp = &pint;"
333 NL " private int * pintp = (private int *)intp;"
334 NL " failures += !(isFenceValid(get_fence(pintp)));"
335 NL " failures += !(to_private(pintp));"
336 NL " failures += (*pintp != 3);"
337 NL
338 NL " results[tid] = (failures == 0);"
339 NL "}"
340 NL
341 );
342
343 CBasicTest test(KERNEL_FUNCTIONS);
344
345 return test.Execute(deviceID, context, queue, num_elements);
346 }
347
test_conditional_casting(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)348 int test_conditional_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
349 const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
350 NL
351 NL "__global int gint = 1;"
352 NL
353 NL "__kernel void testKernel(__global uint *results) {"
354 NL " uint tid = get_global_id(0);"
355 NL
356 NL " int *ptr;"
357 NL " __local int lint;"
358 NL " lint = 2;"
359 NL
360 NL " if (tid % 2)"
361 NL " ptr = &gint;"
362 NL " else"
363 NL " ptr = &lint;"
364 NL
365 NL " barrier(CLK_GLOBAL_MEM_FENCE);"
366 NL
367 NL " if (tid % 2)"
368 NL " results[tid] = (isFenceValid(get_fence(ptr)) && to_global(ptr) && *ptr == 1);"
369 NL " else"
370 NL " results[tid] = (isFenceValid(get_fence(ptr)) && to_local(ptr) && *ptr == 2);"
371 NL "}"
372 NL;
373
374 CBasicTest test(KERNEL_FUNCTION);
375
376 return test.Execute(deviceID, context, queue, num_elements);
377 }
378
test_chain_casting(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)379 int test_chain_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
380 const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
381 NL
382 NL "__global int gint = 1;"
383 NL
384 NL "int f4(int val, int *ptr) { return (isFenceValid(get_fence(ptr)) && val == *ptr) ? 0 : 1; }"
385 NL "int f3(int val, int *ptr) { return f4(val, ptr); }"
386 NL "int f2(int *ptr, int val) { return f3(val, ptr); }"
387 NL "int f1(int *ptr, int val) { return f2(ptr, val); }"
388 NL
389 NL "__kernel void testKernel(__global uint *results) {"
390 NL " uint tid = get_global_id(0);"
391 NL
392 NL " int *ptr;"
393 NL " __local int lint;"
394 NL " lint = 2;"
395 NL " __private int pint = 3;"
396 NL
397 NL " int failures = 0;"
398 NL " failures += f1(&gint, gint);"
399 NL " failures += f1(&lint, lint);"
400 NL " failures += f1(&pint, pint);"
401 NL
402 NL " results[tid] = (failures == 0);"
403 NL "}"
404 NL;
405 CBasicTest test(KERNEL_FUNCTION);
406
407 return test.Execute(deviceID, context, queue, num_elements);
408 }
409
test_ternary_operator_casting(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)410 int test_ternary_operator_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
411 const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE +
412 NL
413 NL "__global int gint = 1;"
414 NL
415 NL "__kernel void testKernel(__global uint *results) {"
416 NL " uint tid = get_global_id(0);"
417 NL
418 NL " int *ptr;"
419 NL " __local int lint;"
420 NL " lint = 2;"
421 NL
422 NL " ptr = (tid % 2) ? &gint : (int *)&lint; // assuming there is an implicit conversion from named address space to generic"
423 NL
424 NL " barrier(CLK_GLOBAL_MEM_FENCE);"
425 NL
426 NL " if (tid % 2)"
427 NL " results[tid] = (isFenceValid(get_fence(ptr)) && to_global(ptr) && *ptr == gint);"
428 NL " else"
429 NL " results[tid] = (isFenceValid(get_fence(ptr)) && to_local(ptr) && *ptr == lint);"
430 NL "}"
431 NL;
432
433 CBasicTest test(KERNEL_FUNCTION);
434
435 return test.Execute(deviceID, context, queue, num_elements);
436 }
437
test_language_struct(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)438 int test_language_struct(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
439 std::vector<std::string> KERNEL_FUNCTIONS;
440
441 // implicit private struct
442 KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
443 NL
444 NL "__global int gint = 1;"
445 NL
446 NL "__kernel void testKernel(__global uint *results) {"
447 NL " uint tid = get_global_id(0);"
448 NL " int failures = 0;"
449 NL
450 NL " __local int lint;"
451 NL " lint = 2;"
452 NL " __private int pint = 3;"
453 NL
454 NL " struct {"
455 NL " __global int *gintp;"
456 NL " __local int *lintp;"
457 NL " __private int *pintp;"
458 NL " } structWithPointers;"
459 NL
460 NL " structWithPointers.gintp = &gint;"
461 NL " structWithPointers.lintp = &lint;"
462 NL " structWithPointers.pintp = &pint;"
463 NL
464 NL " failures += !(isFenceValid(get_fence(structWithPointers.gintp)));"
465 NL " failures += !(isFenceValid(get_fence(structWithPointers.lintp)));"
466 NL " failures += !(isFenceValid(get_fence(structWithPointers.pintp)));"
467 NL
468 NL " failures += !(to_global(structWithPointers.gintp));"
469 NL " failures += !(to_local(structWithPointers.lintp));"
470 NL " failures += !(to_private(structWithPointers.pintp));"
471 NL
472 NL " results[tid] = (failures == 0);"
473 NL "}"
474 NL
475 );
476
477 // explicit __private struct
478 KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
479 NL
480 NL "__global int gint = 1;"
481 NL
482 NL "__kernel void testKernel(__global uint *results) {"
483 NL " uint tid = get_global_id(0);"
484 NL " int failures = 0;"
485 NL
486 NL " __local int lint;"
487 NL " lint = 2;"
488 NL " __private int pint = 3;"
489 NL
490 NL " typedef struct {"
491 NL " __global int * gintp;"
492 NL " __local int * lintp;"
493 NL " __private int * pintp;"
494 NL " } S;"
495 NL
496 NL " __private S structWithPointers;"
497 NL " structWithPointers.gintp = &gint;"
498 NL " structWithPointers.lintp = &lint;"
499 NL " structWithPointers.pintp = &pint;"
500 NL
501 NL " failures += !(isFenceValid(get_fence(structWithPointers.gintp)));"
502 NL " failures += !(isFenceValid(get_fence(structWithPointers.lintp)));"
503 NL " failures += !(isFenceValid(get_fence(structWithPointers.pintp)));"
504 NL
505 NL " failures += !(to_global(structWithPointers.gintp));"
506 NL " failures += !(to_local(structWithPointers.lintp));"
507 NL " failures += !(to_private(structWithPointers.pintp));"
508 NL
509 NL " results[tid] = (failures == 0);"
510 NL "}"
511 NL
512 );
513
514 KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
515 NL
516 NL "__global int gint = 1;"
517 NL
518 NL "__kernel void testKernel(__global uint *results) {"
519 NL " uint tid = get_global_id(0);"
520 NL " int failures = 0;"
521 NL
522 NL " __local int lint;"
523 NL " lint = 2;"
524 NL " __private int pint = 3;"
525 NL
526 NL " typedef struct {"
527 NL " __global int * gintp;"
528 NL " __local int * lintp;"
529 NL " __private int * pintp;"
530 NL " } S;"
531 NL
532 NL " __local S structWithPointers;"
533 NL " structWithPointers.gintp = &gint;"
534 NL " structWithPointers.lintp = &lint;"
535 NL " structWithPointers.pintp = &pint;"
536 NL
537 NL " failures += !(isFenceValid(get_fence(structWithPointers.gintp)));"
538 NL " failures += !(isFenceValid(get_fence(structWithPointers.lintp)));"
539 NL " failures += !(isFenceValid(get_fence(structWithPointers.pintp)));"
540 NL
541 NL " failures += !(to_global(structWithPointers.gintp));"
542 NL " failures += !(to_local(structWithPointers.lintp));"
543 NL " failures += !(to_private(structWithPointers.pintp));"
544 NL
545 NL " results[tid] = (failures == 0);"
546 NL "}"
547 NL
548 );
549
550 KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
551 NL
552 NL "typedef struct {"
553 NL " __global int *gintp;"
554 NL " __local int *lintp;"
555 NL " __private int *pintp;"
556 NL "} S;"
557 NL
558 NL "__global S structWithPointers;"
559 NL "__global int gint = 1;"
560 NL
561 NL "__kernel void testKernel(__global uint *results) {"
562 NL " uint tid = get_global_id(0);"
563 NL " int failures = 0;"
564 NL
565 NL " __local int lint;"
566 NL " lint = 2;"
567 NL " __private int pint = 3;"
568 NL
569 NL " structWithPointers.gintp = &gint;"
570 NL " structWithPointers.lintp = &lint;"
571 NL " structWithPointers.pintp = &pint;"
572 NL
573 NL " failures += !(isFenceValid(get_fence(structWithPointers.gintp)));"
574 NL " failures += !(isFenceValid(get_fence(structWithPointers.lintp)));"
575 NL " failures += !(isFenceValid(get_fence(structWithPointers.pintp)));"
576 NL
577 NL " failures += !(to_global(structWithPointers.gintp));"
578 NL " failures += !(to_local(structWithPointers.lintp));"
579 NL " failures += !(to_private(structWithPointers.pintp));"
580 NL
581 NL " results[tid] = (failures == 0);"
582 NL "}"
583 NL
584 );
585
586 CBasicTest test(KERNEL_FUNCTIONS);
587
588 return test.Execute(deviceID, context, queue, num_elements);
589 }
590
test_language_union(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)591 int test_language_union(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
592 std::vector<std::string> KERNEL_FUNCTIONS;
593
594 KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
595 NL
596 NL "__global int g = 1;"
597 NL
598 NL "__kernel void testKernel(__global uint *results) {"
599 NL " uint tid = get_global_id(0);"
600 NL " int failures = 0;"
601 NL
602 NL " __local int l;"
603 NL " l = 2;"
604 NL " int p = 3;"
605 NL
606 NL " union {"
607 NL " __global int *gintp;"
608 NL " __local int *lintp;"
609 NL " __private int *pintp;"
610 NL " } u;"
611 NL
612 NL " u.gintp = &g;"
613 NL " failures += !(isFenceValid(get_fence(u.gintp)));"
614 NL " failures += !to_global(u.gintp);"
615 NL " failures += (*(u.gintp) != 1);"
616 NL
617 NL " u.lintp = &l;"
618 NL " failures += !(isFenceValid(get_fence(u.lintp)));"
619 NL " failures += !to_local(u.lintp);"
620 NL " failures += (*(u.lintp) != 2);"
621 NL
622 NL " u.pintp = &p;"
623 NL " failures += !(isFenceValid(get_fence(u.pintp)));"
624 NL " failures += !to_private(u.pintp);"
625 NL " failures += (*(u.pintp) != 3);"
626 NL
627 NL " results[tid] = (failures == 0);"
628 NL "}"
629 NL
630 );
631
632 KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
633 NL
634 NL "__global int g = 1;"
635 NL
636 NL "__kernel void testKernel(__global uint *results) {"
637 NL " uint tid = get_global_id(0);"
638 NL " int failures = 0;"
639 NL
640 NL " __local int l;"
641 NL " l = 2;"
642 NL " int p = 3;"
643 NL
644 NL " typedef union {"
645 NL " __global int * gintp;"
646 NL " __local int * lintp;"
647 NL " __private int * pintp;"
648 NL " } U;"
649 NL
650 NL " __local U u;"
651 NL
652 NL " u.gintp = &g;"
653 NL " work_group_barrier(CLK_LOCAL_MEM_FENCE);"
654 NL " failures += !(isFenceValid(get_fence(u.gintp)));"
655 NL " failures += !to_global(u.gintp);"
656 NL " failures += (*(u.gintp) != 1);"
657 NL
658 NL " work_group_barrier(CLK_LOCAL_MEM_FENCE);"
659 NL " u.lintp = &l;"
660 NL " work_group_barrier(CLK_LOCAL_MEM_FENCE);"
661 NL " failures += !(isFenceValid(get_fence(u.lintp)));"
662 NL " failures += !to_local(u.lintp);"
663 NL " failures += (*(u.lintp) != 2);"
664 NL
665 NL " work_group_barrier(CLK_LOCAL_MEM_FENCE);"
666 NL " if(get_local_id(0) == 0) {"
667 NL " u.pintp = &p;"
668 NL " failures += !(isFenceValid(get_fence(u.pintp)));"
669 NL " failures += !to_private(u.pintp);"
670 NL " failures += (*(u.pintp) != 3);"
671 NL " }"
672 NL
673 NL " results[tid] = (failures == 0);"
674 NL "}"
675 NL
676 );
677
678 KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE +
679 NL
680 NL "typedef union {"
681 NL " __global int * gintp;"
682 NL " __local int * lintp;"
683 NL " __private int * pintp;"
684 NL "} U;"
685 NL
686 NL "__global U u;"
687 NL "__global int g = 1;"
688 NL
689 NL "__kernel void testKernel(__global uint *results) {"
690 NL " uint tid = get_global_id(0);"
691 NL
692 NL " // for global unions only one thread should modify union's content"
693 NL " if (tid != 0) {"
694 NL " results[tid] = 1;"
695 NL " return;"
696 NL " }"
697 NL
698 NL " int failures = 0;"
699 NL
700 NL " __local int l;"
701 NL " l = 2;"
702 NL " int p = 3;"
703 NL
704 NL " u.gintp = &g;"
705 NL " failures += !(isFenceValid(get_fence(u.gintp)));"
706 NL " failures += !to_global(u.gintp);"
707 NL " failures += (*(u.gintp) != 1);"
708 NL
709 NL " u.lintp = &l;"
710 NL " failures += !(isFenceValid(get_fence(u.lintp)));"
711 NL " failures += !to_local(u.lintp);"
712 NL " failures += (*(u.lintp) != 2);"
713 NL
714 NL " u.pintp = &p;"
715 NL " failures += !(isFenceValid(get_fence(u.pintp)));"
716 NL " failures += !to_private(u.pintp);"
717 NL " failures += (*(u.pintp) != 3);"
718 NL
719 NL " results[tid] = (failures == 0);"
720 NL "}"
721 NL
722 );
723
724 CBasicTest test(KERNEL_FUNCTIONS);
725
726 return test.Execute(deviceID, context, queue, num_elements);
727 }
728
test_multiple_calls_same_function(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)729 int test_multiple_calls_same_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
730 const std::string KERNEL_FUNCTION =
731 NL
732 NL "int shift2(const int *ptr, int arg) {"
733 NL " return *ptr << arg;"
734 NL "}"
735 NL
736 NL "__kernel void testKernel(__global uint *results) {"
737 NL " uint tid = get_global_id(0);"
738 NL " int failures = 0;"
739 NL
740 NL " __local int val;"
741 NL " val = get_group_id(0);"
742 NL
743 NL " for (int i = 0; i < 5; i++) {"
744 NL " if (shift2(&val, i) != (val << i))"
745 NL " failures++;"
746 NL " }"
747 NL
748 NL " for (int i = 10; i > 5; i--) {"
749 NL " if (shift2(&val, i) != (val << i))"
750 NL " failures++;"
751 NL " }"
752 NL
753 NL " results[tid] = (failures == 0);"
754 NL "}"
755 NL;
756
757 CBasicTest test(KERNEL_FUNCTION);
758
759 return test.Execute(deviceID, context, queue, num_elements);
760 }
761
test_compare_pointers(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)762 int test_compare_pointers(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
763 std::vector<std::string> KERNEL_FUNCTIONS;
764
765 KERNEL_FUNCTIONS.push_back(
766 NL "__kernel void testKernel(__global uint *results) {"
767 NL " uint tid = get_global_id(0);"
768 NL
769 NL " int *ptr = NULL;"
770 NL
771 NL " results[tid] = (ptr == NULL);"
772 NL "}"
773 NL
774 );
775
776 KERNEL_FUNCTIONS.push_back(
777 NL "__kernel void testKernel(__global uint *results) {"
778 NL " uint tid = get_global_id(0);"
779 NL
780 NL " int *ptr = NULL;"
781 NL " __global int *gptr = NULL;"
782 NL
783 NL " results[tid] = (ptr == gptr);"
784 NL "}"
785 NL
786 );
787
788 KERNEL_FUNCTIONS.push_back(
789 NL "__kernel void testKernel(__global uint *results) {"
790 NL " uint tid = get_global_id(0);"
791 NL
792 NL " int *ptr = NULL;"
793 NL " __local int *lptr = NULL;"
794 NL
795 NL " results[tid] = (ptr == lptr);"
796 NL "}"
797 NL
798 );
799
800 KERNEL_FUNCTIONS.push_back(
801 NL "__kernel void testKernel(__global uint *results) {"
802 NL " uint tid = get_global_id(0);"
803 NL
804 NL " int *ptr = NULL;"
805 NL " __private int *pptr = NULL;"
806 NL
807 NL " results[tid] = (ptr == pptr);"
808 NL "}"
809 NL
810 );
811
812 KERNEL_FUNCTIONS.push_back(
813 NL "__kernel void testKernel(__global uint *results) {"
814 NL " uint tid = get_global_id(0);"
815 NL
816 NL " int *ptr = NULL;"
817 NL " __local int *lptr = NULL;"
818 NL " __global int *gptr = NULL;"
819 NL
820 NL " ptr = lptr;"
821 NL
822 NL " results[tid] = (gptr == ptr) && (lptr == ptr);"
823 NL "}"
824 NL
825 );
826
827 KERNEL_FUNCTIONS.push_back(
828 NL "__kernel void testKernel(__global uint *results) {"
829 NL " uint tid = get_global_id(0);"
830 NL
831 NL " int some_value = 7;"
832 NL " int *ptr = NULL;"
833 NL " __private int *pptr = &some_value;"
834 NL
835 NL " results[tid] = (ptr != pptr);"
836 NL "}"
837 NL
838 );
839
840 KERNEL_FUNCTIONS.push_back(
841 NL "__kernel void testKernel(__global uint *results) {"
842 NL " uint tid = get_global_id(0);"
843 NL
844 NL " __local int some_value;"
845 NL " some_value = 7;"
846 NL " int *ptr = NULL;"
847 NL " __local int *lptr = &some_value;"
848 NL
849 NL " results[tid] = (ptr != lptr);"
850 NL "}"
851 NL
852 );
853
854 KERNEL_FUNCTIONS.push_back(
855 NL "__global int some_value = 7;"
856 NL
857 NL "__kernel void testKernel(__global uint *results) {"
858 NL " uint tid = get_global_id(0);"
859 NL
860 NL " int *ptr = NULL;"
861 NL " __global int *gptr = &some_value;"
862 NL
863 NL " results[tid] = (ptr != gptr);"
864 NL "}"
865 NL
866 );
867
868 KERNEL_FUNCTIONS.push_back(
869 NL "__global int arr[5] = { 0, 1, 2, 3, 4 };"
870 NL
871 NL "__kernel void testKernel(__global uint *results) {"
872 NL " uint tid = get_global_id(0);"
873 NL
874 NL " int *ptr = &arr[1];"
875 NL " __global int *gptr = &arr[3];"
876 NL
877 NL " results[tid] = (gptr >= ptr);"
878 NL "}"
879 NL
880 );
881
882 CBasicTest test(KERNEL_FUNCTIONS);
883
884 return test.Execute(deviceID, context, queue, num_elements);
885 }
886