1 //
2 // Copyright (c) 2020 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 "testBase.h"
17 #include "harness/testHarness.h"
18 #include "harness/deviceInfo.h"
19
20 static const char* test_kernel = R"CLC(
21 __kernel void test(__global int* dst) {
22 dst[0] = 0;
23 }
24 )CLC";
25
test_consistency_svm(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)26 int test_consistency_svm(cl_device_id deviceID, cl_context context,
27 cl_command_queue queue, int num_elements)
28 {
29 // clGetDeviceInfo, passing CL_DEVICE_SVM_CAPABILITIES:
30 // May return 0, indicating that device does not support Shared Virtual
31 // Memory.
32 cl_int error;
33
34 const size_t allocSize = 16;
35 clMemWrapper mem;
36 clProgramWrapper program;
37 clKernelWrapper kernel;
38
39 cl_device_svm_capabilities svmCaps = 0;
40 error = clGetDeviceInfo(deviceID, CL_DEVICE_SVM_CAPABILITIES,
41 sizeof(svmCaps), &svmCaps, NULL);
42 test_error(error, "Unable to query CL_DEVICE_SVM_CAPABILITIES");
43
44 if (svmCaps == 0)
45 {
46 // Test setup:
47
48 mem =
49 clCreateBuffer(context, CL_MEM_READ_WRITE, allocSize, NULL, &error);
50 test_error(error, "Unable to create test buffer");
51
52 error = create_single_kernel_helper(context, &program, &kernel, 1,
53 &test_kernel, "test");
54 test_error(error, "Unable to create test kernel");
55
56 // clGetMemObjectInfo, passing CL_MEM_USES_SVM_POINTER
57 // Returns CL_FALSE if no devices in the context associated with
58 // memobj support Shared Virtual Memory.
59 cl_bool usesSVMPointer;
60 error =
61 clGetMemObjectInfo(mem, CL_MEM_USES_SVM_POINTER,
62 sizeof(usesSVMPointer), &usesSVMPointer, NULL);
63 test_error(error, "Unable to query CL_MEM_USES_SVM_POINTER");
64 test_assert_error(usesSVMPointer == CL_FALSE,
65 "CL_DEVICE_SVM_CAPABILITIES returned 0 but "
66 "CL_MEM_USES_SVM_POINTER did not return CL_FALSE");
67
68 // Check that the SVM APIs can be called.
69
70 // Returns NULL if no devices in context support Shared Virtual Memory.
71 void* ptr0 = clSVMAlloc(context, CL_MEM_READ_WRITE, allocSize, 0);
72 void* ptr1 = clSVMAlloc(context, CL_MEM_READ_WRITE, allocSize, 0);
73 test_assert_error(ptr0 == NULL && ptr1 == NULL,
74 "CL_DEVICE_SVM_CAPABILITIES returned 0 but "
75 "clSVMAlloc returned a non-NULL value");
76
77 // clEnqueueSVMFree, clEnqueueSVMMemcpy, clEnqueueSVMMemFill,
78 // clEnqueueSVMMap, clEnqueueSVMUnmap, clEnqueueSVMMigrateMem Returns
79 // CL_INVALID_OPERATION if the device associated with command_queue does
80 // not support Shared Virtual Memory.
81
82 // These calls purposefully pass bogus pointers to the functions to
83 // better test that they are a NOP when SVM is not supported.
84 void* bogus0 = (void*)0xDEADBEEF;
85 void* bogus1 = (void*)0xDEADDEAD;
86 cl_uint pattern = 0xAAAAAAAA;
87 error = clEnqueueSVMMemFill(queue, bogus0, &pattern, sizeof(pattern),
88 allocSize, 0, NULL, NULL);
89 test_failure_error(
90 error, CL_INVALID_OPERATION,
91 "CL_DEVICE_SVM_CAPABILITIES returned 0 but clEnqueueSVMMemFill did "
92 "not return CL_INVALID_OPERATION");
93
94 error = clEnqueueSVMMemcpy(queue, CL_TRUE, bogus1, bogus0, allocSize, 0,
95 NULL, NULL);
96 test_failure_error(
97 error, CL_INVALID_OPERATION,
98 "CL_DEVICE_SVM_CAPABILITIES returned 0 but "
99 "clEnqueueSVMMemcpy did not return CL_INVALID_OPERATION");
100
101 error = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ, bogus1, allocSize,
102 0, NULL, NULL);
103 test_failure_error(
104 error, CL_INVALID_OPERATION,
105 "CL_DEVICE_SVM_CAPABILITIES returned 0 but "
106 "clEnqueueSVMMap did not return CL_INVALID_OPERATION");
107
108 error = clEnqueueSVMUnmap(queue, bogus1, 0, NULL, NULL);
109 test_failure_error(
110 error, CL_INVALID_OPERATION,
111 "CL_DEVICE_SVM_CAPABILITIES returned 0 but "
112 "clEnqueueSVMUnmap did not return CL_INVALID_OPERATION");
113
114 error = clEnqueueSVMMigrateMem(queue, 1, (const void**)&bogus1, NULL, 0,
115 0, NULL, NULL);
116 test_failure_error(
117 error, CL_INVALID_OPERATION,
118 "CL_DEVICE_SVM_CAPABILITIES returned 0 but "
119 "clEnqueueSVMMigrateMem did not return CL_INVALID_OPERATION");
120
121 // If the enqueue calls above did not return errors, a clFinish would be
122 // needed here to ensure the SVM operations are complete before freeing
123 // the SVM pointers.
124
125 clSVMFree(context, bogus0);
126 error = clEnqueueSVMFree(queue, 1, &bogus0, NULL, NULL, 0, NULL, NULL);
127 test_failure_error(
128 error, CL_INVALID_OPERATION,
129 "CL_DEVICE_SVM_CAPABILITIES returned 0 but "
130 "clEnqueueSVMFree did not return CL_INVALID_OPERATION");
131
132 // If the enqueue calls above did not return errors, a clFinish should
133 // be included here to ensure the enqueued SVM free is complete.
134
135 // clSetKernelArgSVMPointer, clSetKernelExecInfo
136 // Returns CL_INVALID_OPERATION if no devices in the context associated
137 // with kernel support Shared Virtual Memory.
138
139 error = clSetKernelArgSVMPointer(kernel, 0, NULL);
140 test_failure_error(
141 error, CL_INVALID_OPERATION,
142 "CL_DEVICE_SVM_CAPABILITIES returned 0 but "
143 "clSetKernelArgSVMPointer did not return CL_INVALID_OPERATION");
144
145 error =
146 clSetKernelExecInfo(kernel, CL_KERNEL_EXEC_INFO_SVM_PTRS, 0, NULL);
147 test_failure_error(
148 error, CL_INVALID_OPERATION,
149 "CL_DEVICE_SVM_CAPABILITIES returned 0 but "
150 "clSetKernelExecInfo did not return CL_INVALID_OPERATION");
151 }
152
153 return TEST_PASS;
154 }
155
check_atomic_capabilities(cl_device_atomic_capabilities atomicCaps,cl_device_atomic_capabilities requiredCaps)156 static int check_atomic_capabilities(cl_device_atomic_capabilities atomicCaps,
157 cl_device_atomic_capabilities requiredCaps)
158 {
159 if ((atomicCaps & requiredCaps) != requiredCaps)
160 {
161 log_error("Atomic capabilities %llx is missing support for at least "
162 "one required capability %llx!\n",
163 atomicCaps, requiredCaps);
164 return TEST_FAIL;
165 }
166
167 if ((atomicCaps & CL_DEVICE_ATOMIC_SCOPE_ALL_DEVICES) != 0
168 && (atomicCaps & CL_DEVICE_ATOMIC_SCOPE_DEVICE) == 0)
169 {
170 log_error("Support for CL_DEVICE_ATOMIC_SCOPE_ALL_DEVICES requires "
171 "support for CL_DEVICE_ATOMIC_SCOPE_DEVICE!\n");
172 return TEST_FAIL;
173 }
174
175 if ((atomicCaps & CL_DEVICE_ATOMIC_SCOPE_DEVICE) != 0
176 && (atomicCaps & CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP) == 0)
177 {
178 log_error("Support for CL_DEVICE_ATOMIC_SCOPE_DEVICE requires "
179 "support for CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP!\n");
180 return TEST_FAIL;
181 }
182
183 if ((atomicCaps & CL_DEVICE_ATOMIC_ORDER_SEQ_CST) != 0
184 && (atomicCaps & CL_DEVICE_ATOMIC_ORDER_ACQ_REL) == 0)
185 {
186 log_error("Support for CL_DEVICE_ATOMIC_ORDER_SEQ_CST requires "
187 "support for CL_DEVICE_ATOMIC_ORDER_ACQ_REL!\n");
188 return TEST_FAIL;
189 }
190
191 if ((atomicCaps & CL_DEVICE_ATOMIC_ORDER_ACQ_REL) != 0
192 && (atomicCaps & CL_DEVICE_ATOMIC_ORDER_RELAXED) == 0)
193 {
194 log_error("Support for CL_DEVICE_ATOMIC_ORDER_ACQ_REL requires "
195 "support for CL_DEVICE_ATOMIC_ORDER_RELAXED!\n");
196 return TEST_FAIL;
197 }
198
199 return TEST_PASS;
200 }
201
test_consistency_memory_model(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)202 int test_consistency_memory_model(cl_device_id deviceID, cl_context context,
203 cl_command_queue queue, int num_elements)
204 {
205 cl_int error;
206 cl_device_atomic_capabilities atomicCaps = 0;
207
208 error = clGetDeviceInfo(deviceID, CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES,
209 sizeof(atomicCaps), &atomicCaps, NULL);
210 test_error(error, "Unable to query CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES");
211
212 error = check_atomic_capabilities(atomicCaps,
213 CL_DEVICE_ATOMIC_ORDER_RELAXED
214 | CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP);
215 if (error == TEST_FAIL)
216 {
217 log_error("Checks failed for CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES\n");
218 return error;
219 }
220
221 error = clGetDeviceInfo(deviceID, CL_DEVICE_ATOMIC_FENCE_CAPABILITIES,
222 sizeof(atomicCaps), &atomicCaps, NULL);
223 test_error(error, "Unable to query CL_DEVICE_ATOMIC_FENCE_CAPABILITIES");
224
225 error = check_atomic_capabilities(atomicCaps,
226 CL_DEVICE_ATOMIC_ORDER_RELAXED
227 | CL_DEVICE_ATOMIC_ORDER_ACQ_REL
228 | CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP);
229 if (error == TEST_FAIL)
230 {
231 log_error("Checks failed for CL_DEVICE_ATOMIC_FENCE_CAPABILITIES\n");
232 return error;
233 }
234
235 return TEST_PASS;
236 }
237
test_consistency_device_enqueue(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)238 int test_consistency_device_enqueue(cl_device_id deviceID, cl_context context,
239 cl_command_queue queue, int num_elements)
240 {
241 // clGetDeviceInfo, passing CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES
242 // May return 0, indicating that device does not support Device-Side Enqueue
243 // and On-Device Queues.
244 cl_int error;
245
246 cl_device_device_enqueue_capabilities dseCaps = 0;
247 error = clGetDeviceInfo(deviceID, CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES,
248 sizeof(dseCaps), &dseCaps, NULL);
249 test_error(error, "Unable to query CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES");
250
251 if (dseCaps == 0)
252 {
253 // clGetDeviceInfo, passing CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES
254 // Returns 0 if device does not support Device-Side Enqueue and
255 // On-Device Queues.
256
257 cl_command_queue_properties devQueueProps = 0;
258 error = clGetDeviceInfo(deviceID, CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES,
259 sizeof(devQueueProps), &devQueueProps, NULL);
260 test_error(error,
261 "Unable to query CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES");
262 test_assert_error(
263 devQueueProps == 0,
264 "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but "
265 "CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES returned a non-zero value");
266
267 // clGetDeviceInfo, passing
268 // CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE,
269 // CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE,
270 // CL_DEVICE_MAX_ON_DEVICE_QUEUES, or
271 // CL_DEVICE_MAX_ON_DEVICE_EVENTS
272 // Returns 0 if device does not support Device-Side Enqueue and
273 // On-Device Queues.
274
275 cl_uint u = 0;
276
277 error =
278 clGetDeviceInfo(deviceID, CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE,
279 sizeof(u), &u, NULL);
280 test_error(error,
281 "Unable to query CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE");
282 test_assert_error(u == 0,
283 "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 "
284 "but CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE "
285 "returned a non-zero value");
286
287 error = clGetDeviceInfo(deviceID, CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE,
288 sizeof(u), &u, NULL);
289 test_error(error, "Unable to query CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE");
290 test_assert_error(
291 u == 0,
292 "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but "
293 "CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE returned a non-zero value");
294
295 error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_ON_DEVICE_QUEUES,
296 sizeof(u), &u, NULL);
297 test_error(error, "Unable to query CL_DEVICE_MAX_ON_DEVICE_QUEUES");
298 test_assert_error(
299 u == 0,
300 "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but "
301 "CL_DEVICE_MAX_ON_DEVICE_QUEUES returned a non-zero value");
302
303 error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_ON_DEVICE_EVENTS,
304 sizeof(u), &u, NULL);
305 test_error(error, "Unable to query CL_DEVICE_MAX_ON_DEVICE_EVENTS");
306 test_assert_error(
307 u == 0,
308 "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but "
309 "CL_DEVICE_MAX_ON_DEVICE_EVENTS returned a non-zero value");
310
311 // clGetCommandQueueInfo, passing CL_QUEUE_SIZE
312 // Returns CL_INVALID_COMMAND_QUEUE since command_queue cannot be a
313 // valid device command-queue.
314
315 error =
316 clGetCommandQueueInfo(queue, CL_QUEUE_SIZE, sizeof(u), &u, NULL);
317 test_failure_error(
318 error, CL_INVALID_COMMAND_QUEUE,
319 "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but "
320 "CL_QUEUE_SIZE did not return CL_INVALID_COMMAND_QUEUE");
321
322 cl_command_queue q = NULL;
323 error = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE_DEFAULT, sizeof(q),
324 &q, NULL);
325 test_error(error, "Unable to query CL_QUEUE_DEVICE_DEFAULT");
326 test_assert_error(
327 q == NULL,
328 "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 but "
329 "CL_QUEUE_DEVICE_DEFAULT returned a non-NULL value");
330
331 // clSetDefaultDeviceCommandQueue
332 // Returns CL_INVALID_OPERATION if device does not support On-Device
333 // Queues.
334 error = clSetDefaultDeviceCommandQueue(context, deviceID, NULL);
335 test_failure_error(error, CL_INVALID_OPERATION,
336 "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES returned 0 "
337 "but clSetDefaultDeviceCommandQueue did not return "
338 "CL_INVALID_OPERATION");
339 }
340 else
341 {
342 if ((dseCaps & CL_DEVICE_QUEUE_REPLACEABLE_DEFAULT) == 0)
343 {
344 // clSetDefaultDeviceCommandQueue
345 // Returns CL_INVALID_OPERATION if device does not support a
346 // replaceable default On-Device Queue.
347 error = clSetDefaultDeviceCommandQueue(context, deviceID, NULL);
348 test_failure_error(
349 error, CL_INVALID_OPERATION,
350 "CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES did not "
351 "include CL_DEVICE_QUEUE_REPLACEABLE_DEFAULT but "
352 "clSetDefaultDeviceCommandQueue did not return "
353 "CL_INVALID_OPERATION");
354 }
355
356 // If CL_DEVICE_QUEUE_REPLACEABLE_DEFAULT is set,
357 // CL_DEVICE_QUEUE_SUPPORTED must also be set.
358 if ((dseCaps & CL_DEVICE_QUEUE_REPLACEABLE_DEFAULT) != 0
359 && (dseCaps & CL_DEVICE_QUEUE_SUPPORTED) == 0)
360 {
361 log_error("DEVICE_QUEUE_REPLACEABLE_DEFAULT is set but "
362 "DEVICE_QUEUE_SUPPORTED is not set\n");
363 return TEST_FAIL;
364 }
365
366 // Devices that set CL_DEVICE_QUEUE_SUPPORTED must also return CL_TRUE
367 // for CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT.
368 if ((dseCaps & CL_DEVICE_QUEUE_SUPPORTED) != 0)
369 {
370 cl_bool b;
371 error = clGetDeviceInfo(deviceID,
372 CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT,
373 sizeof(b), &b, NULL);
374 test_error(
375 error,
376 "Unable to query CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT");
377 test_assert_error(
378 b == CL_TRUE,
379 "DEVICE_QUEUE_SUPPORTED is set but "
380 "CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT returned CL_FALSE");
381 }
382 }
383
384 return TEST_PASS;
385 }
386
test_consistency_pipes(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)387 int test_consistency_pipes(cl_device_id deviceID, cl_context context,
388 cl_command_queue queue, int num_elements)
389 {
390 // clGetDeviceInfo, passing CL_DEVICE_PIPE_SUPPORT
391 // May return CL_FALSE, indicating that device does not support Pipes.
392 cl_int error;
393
394 cl_bool pipeSupport = CL_FALSE;
395 error = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_SUPPORT,
396 sizeof(pipeSupport), &pipeSupport, NULL);
397 test_error(error, "Unable to query CL_DEVICE_PIPE_SUPPORT");
398
399 if (pipeSupport == CL_FALSE)
400 {
401 // clGetDeviceInfo, passing
402 // CL_DEVICE_MAX_PIPE_ARGS,
403 // CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS, or
404 // CL_DEVICE_PIPE_MAX_PACKET_SIZE
405 // Returns 0 if device does not support Pipes.
406
407 cl_uint u = 0;
408
409 error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_PIPE_ARGS, sizeof(u),
410 &u, NULL);
411 test_error(error, "Unable to query CL_DEVICE_MAX_PIPE_ARGS");
412 test_assert_error(u == 0,
413 "CL_DEVICE_PIPE_SUPPORT returned CL_FALSE, but "
414 "CL_DEVICE_MAX_PIPE_ARGS returned a non-zero value");
415
416 error =
417 clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS,
418 sizeof(u), &u, NULL);
419 test_error(error,
420 "Unable to query CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS");
421 test_assert_error(u == 0,
422 "CL_DEVICE_PIPE_SUPPORT returned CL_FALSE, but "
423 "CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS returned "
424 "a non-zero value");
425
426 error = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_MAX_PACKET_SIZE,
427 sizeof(u), &u, NULL);
428 test_error(error, "Unable to query CL_DEVICE_PIPE_MAX_PACKET_SIZE");
429 test_assert_error(
430 u == 0,
431 "CL_DEVICE_PIPE_SUPPORT returned CL_FALSE, but "
432 "CL_DEVICE_PIPE_MAX_PACKET_SIZE returned a non-zero value");
433
434 // clCreatePipe
435 // Returns CL_INVALID_OPERATION if no devices in context support Pipes.
436 clMemWrapper mem =
437 clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, 4, 4, NULL, &error);
438 test_failure_error(error, CL_INVALID_OPERATION,
439 "CL_DEVICE_PIPE_SUPPORT returned CL_FALSE but "
440 "clCreatePipe did not return CL_INVALID_OPERATION");
441
442 // clGetPipeInfo
443 // Returns CL_INVALID_MEM_OBJECT since pipe cannot be a valid pipe
444 // object.
445 clMemWrapper not_a_pipe =
446 clCreateBuffer(context, CL_MEM_READ_WRITE, 4, NULL, &error);
447 test_error(error, "Unable to create non-pipe buffer");
448
449 error =
450 clGetPipeInfo(not_a_pipe, CL_PIPE_PACKET_SIZE, sizeof(u), &u, NULL);
451 test_failure_error(
452 error, CL_INVALID_MEM_OBJECT,
453 "CL_DEVICE_PIPE_SUPPORT returned CL_FALSE but clGetPipeInfo did "
454 "not return CL_INVALID_MEM_OBJECT");
455 }
456 else
457 {
458 // Devices that support pipes must also return CL_TRUE
459 // for CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT.
460 cl_bool b;
461 error =
462 clGetDeviceInfo(deviceID, CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT,
463 sizeof(b), &b, NULL);
464 test_error(error,
465 "Unable to query CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT");
466 test_assert_error(
467 b == CL_TRUE,
468 "CL_DEVICE_PIPE_SUPPORT returned CL_TRUE but "
469 "CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT returned CL_FALSE");
470 }
471
472 return TEST_PASS;
473 }
474
test_consistency_progvar(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)475 int test_consistency_progvar(cl_device_id deviceID, cl_context context,
476 cl_command_queue queue, int num_elements)
477 {
478 // clGetDeviceInfo, passing CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE
479 // May return 0, indicating that device does not support Program Scope
480 // Global Variables.
481 cl_int error;
482
483 clProgramWrapper program;
484 clKernelWrapper kernel;
485
486 size_t maxGlobalVariableSize = 0;
487 error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE,
488 sizeof(maxGlobalVariableSize),
489 &maxGlobalVariableSize, NULL);
490 test_error(error, "Unable to query CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE");
491
492 if (maxGlobalVariableSize == 0)
493 {
494 // Test setup:
495
496 error = create_single_kernel_helper(context, &program, &kernel, 1,
497 &test_kernel, "test");
498 test_error(error, "Unable to create test kernel");
499
500 size_t sz = SIZE_MAX;
501
502 // clGetDeviceInfo, passing
503 // CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE
504 // Returns 0 if device does not support Program Scope Global Variables.
505
506 error = clGetDeviceInfo(deviceID,
507 CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE,
508 sizeof(sz), &sz, NULL);
509 test_error(
510 error,
511 "Unable to query CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE");
512 test_assert_error(
513 sz == 0,
514 "CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE returned 0 but "
515 "CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE returned a "
516 "non-zero value");
517
518 // clGetProgramBuildInfo, passing
519 // CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE
520 // Returns 0 if device does not support Program Scope Global Variables.
521
522 error = clGetProgramBuildInfo(
523 program, deviceID, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE,
524 sizeof(sz), &sz, NULL);
525 test_error(
526 error,
527 "Unable to query CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE");
528 test_assert_error(sz == 0,
529 "CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE returned 0 "
530 "but CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE "
531 "returned a non-zero value");
532 }
533
534 return TEST_PASS;
535 }
536
test_consistency_non_uniform_work_group(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)537 int test_consistency_non_uniform_work_group(cl_device_id deviceID,
538 cl_context context,
539 cl_command_queue queue,
540 int num_elements)
541 {
542 // clGetDeviceInfo, passing CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT:
543 // May return CL_FALSE, indicating that device does not support Non-Uniform
544 // Work Groups.
545 cl_int error;
546
547 const size_t allocSize = 16;
548 clMemWrapper mem;
549 clProgramWrapper program;
550 clKernelWrapper kernel;
551
552 cl_bool nonUniformWorkGroupSupport = CL_FALSE;
553 error = clGetDeviceInfo(deviceID, CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT,
554 sizeof(nonUniformWorkGroupSupport),
555 &nonUniformWorkGroupSupport, NULL);
556 test_error(error,
557 "Unable to query CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT");
558
559 if (nonUniformWorkGroupSupport == CL_FALSE)
560 {
561 // Test setup:
562
563 mem =
564 clCreateBuffer(context, CL_MEM_READ_WRITE, allocSize, NULL, &error);
565 test_error(error, "Unable to create test buffer");
566
567 error = create_single_kernel_helper(context, &program, &kernel, 1,
568 &test_kernel, "test");
569 test_error(error, "Unable to create test kernel");
570
571 error = clSetKernelArg(kernel, 0, sizeof(mem), &mem);
572
573 // clEnqueueNDRangeKernel
574 // Behaves as though Non-Uniform Work Groups were not enabled for
575 // kernel, if the device associated with command_queue does not support
576 // Non-Uniform Work Groups.
577
578 size_t global_work_size[] = { 3, 3, 3 };
579 size_t local_work_size[] = { 2, 2, 2 };
580
581 // First, check that a NULL local work size succeeds.
582 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size,
583 NULL, 0, NULL, NULL);
584 test_error(error,
585 "Unable to enqueue kernel with a NULL local work size");
586
587 error = clFinish(queue);
588 test_error(error, "Error calling clFinish after NULL local work size");
589
590 // 1D non-uniform work group:
591 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size,
592 local_work_size, 0, NULL, NULL);
593 test_failure_error(
594 error, CL_INVALID_WORK_GROUP_SIZE,
595 "CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT returned CL_FALSE but 1D "
596 "clEnqueueNDRangeKernel did not return CL_INVALID_WORK_GROUP_SIZE");
597
598 // 2D non-uniform work group:
599 global_work_size[0] = local_work_size[0];
600 error = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, global_work_size,
601 local_work_size, 0, NULL, NULL);
602 test_failure_error(
603 error, CL_INVALID_WORK_GROUP_SIZE,
604 "CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT returned CL_FALSE but 2D "
605 "clEnqueueNDRangeKernel did not return CL_INVALID_WORK_GROUP_SIZE");
606
607 // 3D non-uniform work group:
608 global_work_size[1] = local_work_size[1];
609 error = clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global_work_size,
610 local_work_size, 0, NULL, NULL);
611 test_failure_error(
612 error, CL_INVALID_WORK_GROUP_SIZE,
613 "CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT returned CL_FALSE but 3D "
614 "clEnqueueNDRangeKernel did not return CL_INVALID_WORK_GROUP_SIZE");
615 }
616
617 return TEST_PASS;
618 }
619
test_consistency_read_write_images(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)620 int test_consistency_read_write_images(cl_device_id deviceID,
621 cl_context context,
622 cl_command_queue queue, int num_elements)
623 {
624 // clGetDeviceInfo, passing
625 // CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS May return 0,
626 // indicating that device does not support Read-Write Images.
627 cl_int error;
628
629 cl_uint maxReadWriteImageArgs = 0;
630 error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS,
631 sizeof(maxReadWriteImageArgs),
632 &maxReadWriteImageArgs, NULL);
633 test_error(error,
634 "Unable to query "
635 "CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS");
636
637 // clGetSupportedImageFormats, passing
638 // CL_MEM_KERNEL_READ_AND_WRITE
639 // Returns an empty set (such as num_image_formats equal to 0), indicating
640 // that no image formats are supported for reading and writing in the same
641 // kernel, if no devices in context support Read-Write Images.
642
643 cl_uint totalReadWriteImageFormats = 0;
644
645 const cl_mem_object_type image_types[] = {
646 CL_MEM_OBJECT_IMAGE1D, CL_MEM_OBJECT_IMAGE1D_BUFFER,
647 CL_MEM_OBJECT_IMAGE2D, CL_MEM_OBJECT_IMAGE3D,
648 CL_MEM_OBJECT_IMAGE1D_ARRAY, CL_MEM_OBJECT_IMAGE2D_ARRAY,
649 };
650 for (int i = 0; i < ARRAY_SIZE(image_types); i++)
651 {
652 cl_uint numImageFormats = 0;
653 error = clGetSupportedImageFormats(
654 context, CL_MEM_KERNEL_READ_AND_WRITE, image_types[i], 0, NULL,
655 &numImageFormats);
656 test_error(error,
657 "Unable to query number of CL_MEM_KERNEL_READ_AND_WRITE "
658 "image formats");
659
660 totalReadWriteImageFormats += numImageFormats;
661 }
662
663 if (maxReadWriteImageArgs == 0)
664 {
665 test_assert_error(
666 totalReadWriteImageFormats == 0,
667 "CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS returned 0 "
668 "but clGetSupportedImageFormats(CL_MEM_KERNEL_READ_AND_WRITE) "
669 "returned a non-empty set");
670 }
671 else
672 {
673 test_assert_error(
674 totalReadWriteImageFormats != 0,
675 "CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS is non-zero "
676 "but clGetSupportedImageFormats(CL_MEM_KERNEL_READ_AND_WRITE) "
677 "returned an empty set");
678 }
679
680 return TEST_PASS;
681 }
682
test_consistency_2d_image_from_buffer(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)683 int test_consistency_2d_image_from_buffer(cl_device_id deviceID,
684 cl_context context,
685 cl_command_queue queue,
686 int num_elements)
687 {
688 // clGetDeviceInfo, passing CL_DEVICE_IMAGE_PITCH_ALIGNMENT or
689 // CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT
690 // May return 0, indicating that device does not support Creating a 2D Image
691 // from a Buffer.
692 cl_int error;
693
694 const cl_image_format imageFormat = { CL_RGBA, CL_UNORM_INT8 };
695 const size_t imageDim = 2;
696 const size_t elementSize = 4;
697 const size_t bufferSize = imageDim * imageDim * elementSize;
698
699 clMemWrapper buffer;
700 clMemWrapper image;
701
702 cl_uint imagePitchAlignment = 0;
703 error = clGetDeviceInfo(deviceID, CL_DEVICE_IMAGE_PITCH_ALIGNMENT,
704 sizeof(imagePitchAlignment), &imagePitchAlignment,
705 NULL);
706 test_error(error,
707 "Unable to query "
708 "CL_DEVICE_IMAGE_PITCH_ALIGNMENT");
709
710 cl_uint imageBaseAddressAlignment = 0;
711 error = clGetDeviceInfo(deviceID, CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT,
712 sizeof(imageBaseAddressAlignment),
713 &imageBaseAddressAlignment, NULL);
714 test_error(error,
715 "Unable to query "
716 "CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT");
717
718 bool supports_cl_khr_image2d_from_buffer =
719 is_extension_available(deviceID, "cl_khr_image2d_from_buffer");
720
721 if (imagePitchAlignment == 0 || imageBaseAddressAlignment == 0)
722 {
723 // This probably means that Creating a 2D Image from a Buffer is not
724 // supported.
725
726 // Test setup:
727 buffer =
728 clCreateBuffer(context, CL_MEM_READ_ONLY, bufferSize, NULL, &error);
729 test_error(error, "Unable to create test buffer");
730
731 // Check that both queries return zero:
732 test_assert_error(
733 imagePitchAlignment == 0,
734 "CL_DEVICE_IMAGE_PITCH_ALIGNMENT returned a non-zero value but "
735 "CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT returned 0");
736 test_assert_error(
737 imageBaseAddressAlignment == 0,
738 "CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT returned a non-zero value "
739 "but CL_DEVICE_IMAGE_PITCH_ALIGNMENT returned 0");
740
741 // clGetDeviceInfo, passing CL_DEVICE_EXTENSIONS
742 // Will not describe support for the cl_khr_image2d_from_buffer
743 // extension if device does not support Creating a 2D Image from a
744 // Buffer.
745 test_assert_error(supports_cl_khr_image2d_from_buffer == false,
746 "Device does not support Creating a 2D Image from a "
747 "Buffer but does support cl_khr_image2d_from_buffer");
748
749 // clCreateImage or clCreateImageWithProperties, passing image_type
750 // equal to CL_MEM_OBJECT_IMAGE2D and mem_object not equal to
751 // NULL
752 // Returns CL_INVALID_OPERATION if no devices in context support
753 // Creating a 2D Image from a Buffer.
754
755 cl_image_desc imageDesc = { 0 };
756 imageDesc.image_type = CL_MEM_OBJECT_IMAGE2D;
757 imageDesc.image_width = imageDim;
758 imageDesc.image_height = imageDim;
759 imageDesc.mem_object = buffer;
760
761 image = clCreateImage(context, CL_MEM_READ_ONLY, &imageFormat,
762 &imageDesc, NULL, &error);
763 test_failure_error(
764 error, CL_INVALID_OPERATION,
765 "Device does not support Creating a 2D Image from a "
766 "Buffer but clCreateImage did not return CL_INVALID_OPERATION");
767
768 image =
769 clCreateImageWithProperties(context, NULL, CL_MEM_READ_ONLY,
770 &imageFormat, &imageDesc, NULL, &error);
771 test_failure_error(error, CL_INVALID_OPERATION,
772 "Device does not support Creating a 2D Image from a "
773 "Buffer but clCreateImageWithProperties did not "
774 "return CL_INVALID_OPERATION");
775 }
776 else
777 {
778 test_assert_error(supports_cl_khr_image2d_from_buffer,
779 "Device supports Creating a 2D Image from a Buffer "
780 "but does not support cl_khr_image2d_from_buffer");
781 }
782
783 return TEST_PASS;
784 }
785
786 // Nothing needed for sRGB Images:
787 // All of the sRGB Image Channel Orders (such as CL_sRGBA) are optional for
788 // devices supporting OpenCL 3.0.
789
test_consistency_depth_images(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)790 int test_consistency_depth_images(cl_device_id deviceID, cl_context context,
791 cl_command_queue queue, int num_elements)
792 {
793 // The CL_DEPTH Image Channel Order is optional for devices supporting
794 // OpenCL 3.0.
795 cl_int error;
796
797 cl_uint totalDepthImageFormats = 0;
798
799 const cl_mem_flags mem_flags[] = {
800 CL_MEM_WRITE_ONLY,
801 CL_MEM_READ_WRITE,
802 CL_MEM_KERNEL_READ_AND_WRITE,
803 };
804 for (int i = 0; i < ARRAY_SIZE(mem_flags); i++)
805 {
806 cl_uint numImageFormats = 0;
807 error = clGetSupportedImageFormats(context, mem_flags[i],
808 CL_MEM_OBJECT_IMAGE2D, 0, NULL,
809 &numImageFormats);
810 test_error(
811 error,
812 "Unable to query number of CL_MEM_OBJECT_IMAGE2D image formats");
813
814 std::vector<cl_image_format> imageFormats(numImageFormats);
815 error = clGetSupportedImageFormats(
816 context, mem_flags[i], CL_MEM_OBJECT_IMAGE2D, imageFormats.size(),
817 imageFormats.data(), NULL);
818 test_error(error,
819 "Unable to query CL_MEM_OBJECT_IMAGE2D image formats");
820 for (auto& imageFormat : imageFormats)
821 {
822 if (imageFormat.image_channel_order == CL_DEPTH)
823 {
824 totalDepthImageFormats++;
825 }
826 }
827 }
828
829 bool supports_cl_khr_depth_images =
830 is_extension_available(deviceID, "cl_khr_depth_images");
831
832 if (totalDepthImageFormats == 0)
833 {
834 test_assert_error(supports_cl_khr_depth_images == false,
835 "Device does not support Depth Images but does "
836 "support cl_khr_depth_images");
837 }
838 else
839 {
840 test_assert_error(supports_cl_khr_depth_images,
841 "Device supports Depth Images but does not support "
842 "cl_khr_depth_images");
843 }
844
845 return TEST_PASS;
846 }
847
test_consistency_device_and_host_timer(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)848 int test_consistency_device_and_host_timer(cl_device_id deviceID,
849 cl_context context,
850 cl_command_queue queue,
851 int num_elements)
852 {
853 // clGetPlatformInfo, passing CL_PLATFORM_HOST_TIMER_RESOLUTION
854 // May return 0, indicating that platform does not support Device and Host
855 // Timer Synchronization.
856 cl_int error;
857
858 cl_platform_id platform = NULL;
859 error = clGetDeviceInfo(deviceID, CL_DEVICE_PLATFORM, sizeof(platform),
860 &platform, NULL);
861 test_error(error, "Unable to query CL_DEVICE_PLATFORM");
862
863 cl_ulong hostTimerResolution = 0;
864 error = clGetPlatformInfo(platform, CL_PLATFORM_HOST_TIMER_RESOLUTION,
865 sizeof(hostTimerResolution), &hostTimerResolution,
866 NULL);
867 test_error(error, "Unable to query CL_PLATFORM_HOST_TIMER_RESOLUTION");
868
869 if (hostTimerResolution == 0)
870 {
871 // clGetDeviceAndHostTimer, clGetHostTimer
872 // Returns CL_INVALID_OPERATION if the platform associated with device
873 // does not support Device and Host Timer Synchronization.
874
875 cl_ulong dt = 0;
876 cl_ulong ht = 0;
877
878 error = clGetDeviceAndHostTimer(deviceID, &dt, &ht);
879 test_failure_error(
880 error, CL_INVALID_OPERATION,
881 "CL_PLATFORM_HOST_TIMER_RESOLUTION returned 0 but "
882 "clGetDeviceAndHostTimer did not return CL_INVALID_OPERATION");
883
884 error = clGetHostTimer(deviceID, &ht);
885 test_failure_error(
886 error, CL_INVALID_OPERATION,
887 "CL_PLATFORM_HOST_TIMER_RESOLUTION returned 0 but "
888 "clGetHostTimer did not return CL_INVALID_OPERATION");
889 }
890
891 return TEST_PASS;
892 }
893
test_consistency_il_programs(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)894 int test_consistency_il_programs(cl_device_id deviceID, cl_context context,
895 cl_command_queue queue, int num_elements)
896 {
897 // clGetDeviceInfo, passing CL_DEVICE_IL_VERSION or
898 // CL_DEVICE_ILS_WITH_VERSION
899 // May return an empty string and empty array, indicating that device does
900 // not support Intermediate Language Programs.
901 cl_int error;
902
903 clProgramWrapper program;
904 clKernelWrapper kernel;
905
906 // Even if the device does not support Intermediate Language Programs the
907 // size of the string query should not be zero.
908 size_t sz = SIZE_MAX;
909 error = clGetDeviceInfo(deviceID, CL_DEVICE_IL_VERSION, 0, NULL, &sz);
910 test_error(error, "Unable to query CL_DEVICE_IL_VERSION");
911 test_assert_error(sz != 0,
912 "CL_DEVICE_IL_VERSION should return a non-zero size");
913
914 std::string ilVersion = get_device_il_version_string(deviceID);
915
916 error = clGetDeviceInfo(deviceID, CL_DEVICE_ILS_WITH_VERSION, 0, NULL, &sz);
917 test_error(error, "Unable to query CL_DEVICE_ILS_WITH_VERSION");
918
919 if (ilVersion == "" || sz == 0)
920 {
921 // This probably means that Intermediate Language Programs are not
922 // supported.
923
924 // Check that both queries are consistent:
925 test_assert_error(
926 ilVersion == "",
927 "CL_DEVICE_IL_VERSION returned a non-empty string but "
928 "CL_DEVICE_ILS_WITH_VERSION returned no supported ILs");
929
930 test_assert_error(sz == 0,
931 "CL_DEVICE_ILS_WITH_VERSION returned supported ILs "
932 "but CL_DEVICE_IL_VERSION returned an empty string");
933
934 bool supports_cl_khr_il_program =
935 is_extension_available(deviceID, "cl_khr_il_program");
936 test_assert_error(supports_cl_khr_il_program == false,
937 "Device does not support IL Programs but does "
938 "support cl_khr_il_program");
939
940 // Test setup:
941
942 error = create_single_kernel_helper(context, &program, &kernel, 1,
943 &test_kernel, "test");
944 test_error(error, "Unable to create test kernel");
945
946 // clGetProgramInfo, passing CL_PROGRAM_IL
947 // Returns an empty buffer (such as param_value_size_ret equal to 0) if
948 // no devices in the context associated with program support
949 // Intermediate Language Programs.
950
951 error = clGetProgramInfo(program, CL_PROGRAM_IL, 0, NULL, &sz);
952 test_error(error, "Unable to query CL_PROGRAM_IL");
953 test_assert_error(sz == 0,
954 "Device does not support IL Programs but "
955 "CL_PROGRAM_IL returned a non-zero size");
956
957 // clCreateProgramWithIL
958 // Returns CL_INVALID_OPERATION if no devices in context support
959 // Intermediate Language Programs.
960
961 cl_uint bogus = 0xDEADBEEF;
962 clProgramWrapper ilProgram =
963 clCreateProgramWithIL(context, &bogus, sizeof(bogus), &error);
964 test_failure_error(
965 error, CL_INVALID_OPERATION,
966 "Device does not support IL Programs but clCreateProgramWithIL did "
967 "not return CL_INVALID_OPERATION");
968
969 // clSetProgramSpecializationConstant
970 // Returns CL_INVALID_OPERATION if no devices associated with program
971 // support Intermediate Language Programs.
972
973 cl_uint specConst = 42;
974 error = clSetProgramSpecializationConstant(
975 program, 0, sizeof(specConst), &specConst);
976 test_failure_error(error, CL_INVALID_OPERATION,
977 "Device does not support IL Programs but "
978 "clSetProgramSpecializationConstant did not return "
979 "CL_INVALID_OPERATION");
980 }
981
982 return TEST_PASS;
983 }
984
test_consistency_subgroups(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)985 int test_consistency_subgroups(cl_device_id deviceID, cl_context context,
986 cl_command_queue queue, int num_elements)
987 {
988 // clGetDeviceInfo, passing CL_DEVICE_MAX_NUM_SUB_GROUPS
989 // May return 0, indicating that device does not support Subgroups.
990 cl_int error;
991
992 clProgramWrapper program;
993 clKernelWrapper kernel;
994
995 cl_uint maxNumSubGroups = 0;
996 error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_NUM_SUB_GROUPS,
997 sizeof(maxNumSubGroups), &maxNumSubGroups, NULL);
998 test_error(error, "Unable to query CL_DEVICE_MAX_NUM_SUB_GROUPS");
999
1000 if (maxNumSubGroups == 0)
1001 {
1002 // Test setup:
1003
1004 error = create_single_kernel_helper(context, &program, &kernel, 1,
1005 &test_kernel, "test");
1006 test_error(error, "Unable to create test kernel");
1007
1008 // clGetDeviceInfo, passing
1009 // CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS
1010 // Returns CL_FALSE if device does not support Subgroups.
1011
1012 cl_bool ifp = CL_FALSE;
1013 error = clGetDeviceInfo(
1014 deviceID, CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS,
1015 sizeof(ifp), &ifp, NULL);
1016 test_error(
1017 error,
1018 "Unable to query CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS");
1019 test_assert_error(ifp == CL_FALSE,
1020 "Device does not support Subgroups but "
1021 "CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS "
1022 "did not return CL_FALSE");
1023
1024 // clGetDeviceInfo, passing CL_DEVICE_EXTENSIONS
1025 // Will not describe support for the cl_khr_subgroups extension if
1026 // device does not support Subgroups.
1027
1028 bool supports_cl_khr_subgroups =
1029 is_extension_available(deviceID, "cl_khr_subgroups");
1030 test_assert_error(supports_cl_khr_subgroups == false,
1031 "Device does not support Subgroups but does "
1032 "support cl_khr_subgroups");
1033
1034 // clGetKernelSubGroupInfo
1035 // Returns CL_INVALID_OPERATION if device does not support Subgroups.
1036
1037 size_t sz = SIZE_MAX;
1038 error = clGetKernelSubGroupInfo(kernel, deviceID,
1039 CL_KERNEL_MAX_NUM_SUB_GROUPS, 0, NULL,
1040 sizeof(sz), &sz, NULL);
1041 test_failure_error(
1042 error, CL_INVALID_OPERATION,
1043 "Device does not support Subgroups but clGetKernelSubGroupInfo did "
1044 "not return CL_INVALID_OPERATION");
1045 }
1046
1047 return TEST_PASS;
1048 }
1049
program_callback(cl_program,void *)1050 static void CL_CALLBACK program_callback(cl_program, void*) {}
1051
test_consistency_prog_ctor_dtor(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1052 int test_consistency_prog_ctor_dtor(cl_device_id deviceID, cl_context context,
1053 cl_command_queue queue, int num_elements)
1054 {
1055 cl_int error;
1056
1057 clProgramWrapper program;
1058 clKernelWrapper kernel;
1059
1060 // Test setup:
1061
1062 error = create_single_kernel_helper(context, &program, &kernel, 1,
1063 &test_kernel, "test");
1064 test_error(error, "Unable to create test kernel");
1065
1066 // clGetProgramInfo, passing CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT or
1067 // CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT
1068 // Returns CL_FALSE if no devices in the context associated with program
1069 // support Program Initialization and Clean-Up Kernels.
1070
1071 cl_bool b = CL_FALSE;
1072
1073 error = clGetProgramInfo(program, CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT,
1074 sizeof(b), &b, NULL);
1075 test_error(error, "Unable to query CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT");
1076 test_assert_error(
1077 b == CL_FALSE,
1078 "CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT did not return CL_FALSE");
1079
1080 error = clGetProgramInfo(program, CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT,
1081 sizeof(b), &b, NULL);
1082 test_error(error, "Unable to query CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT");
1083 test_assert_error(
1084 b == CL_FALSE,
1085 "CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT did not return CL_FALSE");
1086
1087 // clSetProgramReleaseCallback
1088 // Returns CL_INVALID_OPERATION if no devices in the context associated with
1089 // program support Program Initialization and Clean-Up Kernels.
1090
1091 error = clSetProgramReleaseCallback(program, program_callback, NULL);
1092 test_failure_error(
1093 error, CL_INVALID_OPERATION,
1094 "clSetProgramReleaseCallback did not return CL_INVALID_OPERATION");
1095
1096 return TEST_PASS;
1097 }
1098
test_consistency_3d_image_writes(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1099 int test_consistency_3d_image_writes(cl_device_id deviceID, cl_context context,
1100 cl_command_queue queue, int num_elements)
1101 {
1102 // clGetSupportedImageFormats, passing CL_MEM_OBJECT_IMAGE3D and one of
1103 // CL_MEM_WRITE_ONLY, CL_MEM_READ_WRITE, or CL_MEM_KERNEL_READ_AND_WRITE
1104 // Returns an empty set (such as num_image_formats equal to 0),
1105 // indicating that no image formats are supported for writing to 3D
1106 // image objects, if no devices in context support Writing to 3D Image
1107 // Objects.
1108 cl_int error;
1109
1110 cl_uint total3DImageWriteFormats = 0;
1111
1112 const cl_mem_flags mem_flags[] = {
1113 CL_MEM_WRITE_ONLY,
1114 CL_MEM_READ_WRITE,
1115 CL_MEM_KERNEL_READ_AND_WRITE,
1116 };
1117 for (int i = 0; i < ARRAY_SIZE(mem_flags); i++)
1118 {
1119 cl_uint numImageFormats = 0;
1120 error = clGetSupportedImageFormats(context, mem_flags[i],
1121 CL_MEM_OBJECT_IMAGE3D, 0, NULL,
1122 &numImageFormats);
1123 test_error(
1124 error,
1125 "Unable to query number of CL_MEM_OBJECT_IMAGE3D image formats");
1126
1127 total3DImageWriteFormats += numImageFormats;
1128 }
1129
1130 bool supports_cl_khr_3d_image_writes =
1131 is_extension_available(deviceID, "cl_khr_3d_image_writes");
1132
1133 if (total3DImageWriteFormats == 0)
1134 {
1135 // clGetDeviceInfo, passing CL_DEVICE_EXTENSIONS
1136 // Will not describe support for the cl_khr_3d_image_writes extension if
1137 // device does not support Writing to 3D Image Objects.
1138 test_assert_error(supports_cl_khr_3d_image_writes == false,
1139 "Device does not support Writing to 3D Image Objects "
1140 "but does support cl_khr_3d_image_writes");
1141 }
1142 else
1143 {
1144 test_assert_error(supports_cl_khr_3d_image_writes,
1145 "Device supports Writing to 3D Image Objects but "
1146 "does not support cl_khr_3d_image_writes");
1147 }
1148
1149 return TEST_PASS;
1150 }
1151