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 <vector>
18 #include <algorithm>
19 #include "errorHelpers.h"
20
21 const char* macro_supported_source = R"(kernel void enabled(global int * buf) {
22 int n = get_global_id(0);
23 buf[n] = 0;
24 #ifndef %s
25 #error Feature macro was not defined
26 #endif
27 })";
28
29 const char* macro_not_supported_source =
30 R"(kernel void not_enabled(global int * buf) {
31 int n = get_global_id(0);
32 buf[n] = 0;
33 #ifdef %s
34 #error Feature macro was defined
35 #endif
36 })";
37
38 template <typename T>
check_api_feature_info_capabilities(cl_device_id deviceID,cl_context context,cl_bool & status,cl_device_info check_property,cl_bitfield check_cap)39 cl_int check_api_feature_info_capabilities(cl_device_id deviceID,
40 cl_context context, cl_bool& status,
41 cl_device_info check_property,
42 cl_bitfield check_cap)
43 {
44 cl_int error = CL_SUCCESS;
45 T response;
46 error = clGetDeviceInfo(deviceID, check_property, sizeof(response),
47 &response, NULL);
48 test_error(error, "clGetDeviceInfo failed.\n");
49
50 if ((response & check_cap) == check_cap)
51 {
52 status = CL_TRUE;
53 }
54 else
55 {
56 status = CL_FALSE;
57 }
58 return error;
59 }
60
check_api_feature_info_support(cl_device_id deviceID,cl_context context,cl_bool & status,cl_device_info check_property)61 cl_int check_api_feature_info_support(cl_device_id deviceID, cl_context context,
62 cl_bool& status,
63 cl_device_info check_property)
64 {
65 cl_int error = CL_SUCCESS;
66 cl_bool response;
67 error = clGetDeviceInfo(deviceID, check_property, sizeof(response),
68 &response, NULL);
69 test_error(error, "clGetDeviceInfo failed.\n");
70 status = response;
71 return error;
72 }
73
74 template <typename T>
check_api_feature_info_number(cl_device_id deviceID,cl_context context,cl_bool & status,cl_device_info check_property)75 cl_int check_api_feature_info_number(cl_device_id deviceID, cl_context context,
76 cl_bool& status,
77 cl_device_info check_property)
78 {
79 cl_int error = CL_SUCCESS;
80 T response;
81 error = clGetDeviceInfo(deviceID, check_property, sizeof(response),
82 &response, NULL);
83 test_error(error, "clGetDeviceInfo failed.\n");
84 if (response > 0)
85 {
86 status = CL_TRUE;
87 }
88 else
89 {
90 status = CL_FALSE;
91 }
92 return error;
93 }
94
check_api_feature_info_supported_image_formats(cl_device_id deviceID,cl_context context,cl_bool & status)95 cl_int check_api_feature_info_supported_image_formats(cl_device_id deviceID,
96 cl_context context,
97 cl_bool& status)
98 {
99 cl_int error = CL_SUCCESS;
100 cl_uint response = 0;
101 cl_uint image_format_count;
102 error = clGetSupportedImageFormats(context, CL_MEM_WRITE_ONLY,
103 CL_MEM_OBJECT_IMAGE3D, 0, NULL,
104 &image_format_count);
105 test_error(error, "clGetSupportedImageFormats failed");
106 response += image_format_count;
107 error = clGetSupportedImageFormats(context, CL_MEM_READ_WRITE,
108 CL_MEM_OBJECT_IMAGE3D, 0, NULL,
109 &image_format_count);
110 test_error(error, "clGetSupportedImageFormats failed");
111 response += image_format_count;
112 error = clGetSupportedImageFormats(context, CL_MEM_KERNEL_READ_AND_WRITE,
113 CL_MEM_OBJECT_IMAGE3D, 0, NULL,
114 &image_format_count);
115 test_error(error, "clGetSupportedImageFormats failed");
116 response += image_format_count;
117 if (response > 0)
118 {
119 status = CL_TRUE;
120 }
121 else
122 {
123 status = CL_FALSE;
124 }
125 return error;
126 }
127
check_compiler_feature_info(cl_device_id deviceID,cl_context context,std::string feature_macro,cl_bool & status)128 cl_int check_compiler_feature_info(cl_device_id deviceID, cl_context context,
129 std::string feature_macro, cl_bool& status)
130 {
131 cl_int error = CL_SUCCESS;
132 clProgramWrapper program_supported;
133 clProgramWrapper program_not_supported;
134 char kernel_supported_src[1024];
135 char kernel_not_supported_src[1024];
136 sprintf(kernel_supported_src, macro_supported_source,
137 feature_macro.c_str());
138 const char* ptr_supported = kernel_supported_src;
139 const char* build_options = "-cl-std=CL3.0";
140
141 error = create_single_kernel_helper_create_program(
142 context, &program_supported, 1, &ptr_supported, build_options);
143 test_error(error, "create_single_kernel_helper_create_program failed.\n");
144
145 sprintf(kernel_not_supported_src, macro_not_supported_source,
146 feature_macro.c_str());
147 const char* ptr_not_supported = kernel_not_supported_src;
148 error = create_single_kernel_helper_create_program(
149 context, &program_not_supported, 1, &ptr_not_supported,
150 "-cl-std=CL3.0");
151 test_error(error, "create_single_kernel_helper_create_program failed.\n");
152
153 cl_int status_supported = CL_SUCCESS;
154 cl_int status_not_supported = CL_SUCCESS;
155 status_supported = clBuildProgram(program_supported, 1, &deviceID,
156 build_options, NULL, NULL);
157 status_not_supported = clBuildProgram(program_not_supported, 1, &deviceID,
158 build_options, NULL, NULL);
159 if (status_supported != status_not_supported)
160 {
161 if (status_not_supported == CL_SUCCESS)
162 {
163 // kernel which verifies not supporting return passed
164 status = CL_FALSE;
165 }
166 else
167 {
168 // kernel which verifies supporting return passed
169 status = CL_TRUE;
170 }
171 }
172 else
173 {
174 log_error("Error: The macro feature is defined and undefined "
175 "in the same time\n");
176 error = OutputBuildLogs(program_supported, 1, &deviceID);
177 test_error(error, "OutputBuildLogs failed.\n");
178 error = OutputBuildLogs(program_not_supported, 1, &deviceID);
179 test_error(error, "OutputBuildLogs failed.\n");
180 return TEST_FAIL;
181 }
182 return error;
183 }
184
feature_macro_verify_results(std::string test_macro_name,cl_bool api_status,cl_bool compiler_status,cl_bool & supported)185 int feature_macro_verify_results(std::string test_macro_name,
186 cl_bool api_status, cl_bool compiler_status,
187 cl_bool& supported)
188 {
189 cl_int error = TEST_PASS;
190 log_info("Feature status: API - %s, compiler - %s\n",
191 api_status == CL_TRUE ? "supported" : "not supported",
192 compiler_status == CL_TRUE ? "supported" : "not supported");
193 if (api_status != compiler_status)
194 {
195 log_info("%s - failed\n", test_macro_name.c_str());
196 supported = CL_FALSE;
197 return TEST_FAIL;
198 }
199 else
200 {
201 log_info("%s - passed\n", test_macro_name.c_str());
202 }
203 supported = api_status;
204 return error;
205 }
206
test_feature_macro_atomic_order_acq_rel(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)207 int test_feature_macro_atomic_order_acq_rel(cl_device_id deviceID,
208 cl_context context,
209 std::string test_macro_name,
210 cl_bool& supported)
211 {
212 cl_int error = TEST_FAIL;
213 cl_bool api_status;
214 cl_bool compiler_status;
215 log_info("\n%s ...\n", test_macro_name.c_str());
216 error = check_api_feature_info_capabilities<cl_device_atomic_capabilities>(
217 deviceID, context, api_status, CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES,
218 CL_DEVICE_ATOMIC_ORDER_ACQ_REL);
219 if (error != CL_SUCCESS)
220 {
221 return error;
222 }
223
224 error = check_compiler_feature_info(deviceID, context, test_macro_name,
225 compiler_status);
226 if (error != CL_SUCCESS)
227 {
228 return error;
229 }
230
231 return feature_macro_verify_results(test_macro_name, api_status,
232 compiler_status, supported);
233 }
234
test_feature_macro_atomic_order_seq_cst(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)235 int test_feature_macro_atomic_order_seq_cst(cl_device_id deviceID,
236 cl_context context,
237 std::string test_macro_name,
238 cl_bool& supported)
239 {
240 cl_int error = TEST_FAIL;
241 cl_bool api_status;
242 cl_bool compiler_status;
243 log_info("\n%s ...\n", test_macro_name.c_str());
244
245 error = check_api_feature_info_capabilities<cl_device_atomic_capabilities>(
246 deviceID, context, api_status, CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES,
247 CL_DEVICE_ATOMIC_ORDER_SEQ_CST);
248 if (error != CL_SUCCESS)
249 {
250 return error;
251 }
252
253 error = check_compiler_feature_info(deviceID, context, test_macro_name,
254 compiler_status);
255 if (error != CL_SUCCESS)
256 {
257 return error;
258 }
259
260 return feature_macro_verify_results(test_macro_name, api_status,
261 compiler_status, supported);
262 }
263
test_feature_macro_atomic_scope_device(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)264 int test_feature_macro_atomic_scope_device(cl_device_id deviceID,
265 cl_context context,
266 std::string test_macro_name,
267 cl_bool& supported)
268 {
269 cl_int error = TEST_FAIL;
270 cl_bool api_status;
271 cl_bool compiler_status;
272 log_info("\n%s ...\n", test_macro_name.c_str());
273 error = check_api_feature_info_capabilities<cl_device_atomic_capabilities>(
274 deviceID, context, api_status, CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES,
275 CL_DEVICE_ATOMIC_SCOPE_DEVICE);
276 if (error != CL_SUCCESS)
277 {
278 return error;
279 }
280 error = check_compiler_feature_info(deviceID, context, test_macro_name,
281 compiler_status);
282 if (error != CL_SUCCESS)
283 {
284 return error;
285 }
286
287 return feature_macro_verify_results(test_macro_name, api_status,
288 compiler_status, supported);
289 }
290
test_feature_macro_atomic_scope_all_devices(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)291 int test_feature_macro_atomic_scope_all_devices(cl_device_id deviceID,
292 cl_context context,
293 std::string test_macro_name,
294 cl_bool& supported)
295 {
296 cl_int error = TEST_FAIL;
297 cl_bool api_status;
298 cl_bool compiler_status;
299 log_info("\n%s ...\n", test_macro_name.c_str());
300 error = check_api_feature_info_capabilities<cl_device_atomic_capabilities>(
301 deviceID, context, api_status, CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES,
302 CL_DEVICE_ATOMIC_SCOPE_ALL_DEVICES);
303 if (error != CL_SUCCESS)
304 {
305 return error;
306 }
307 error = check_compiler_feature_info(deviceID, context, test_macro_name,
308 compiler_status);
309 if (error != CL_SUCCESS)
310 {
311 return error;
312 }
313
314 return feature_macro_verify_results(test_macro_name, api_status,
315 compiler_status, supported);
316 }
317
test_feature_macro_3d_image_writes(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)318 int test_feature_macro_3d_image_writes(cl_device_id deviceID,
319 cl_context context,
320 std::string test_macro_name,
321 cl_bool& supported)
322 {
323 cl_int error = TEST_FAIL;
324 cl_bool api_status;
325 cl_bool compiler_status;
326 log_info("\n%s ...\n", test_macro_name.c_str());
327 error = check_api_feature_info_supported_image_formats(deviceID, context,
328 api_status);
329 if (error != CL_SUCCESS)
330 {
331 return error;
332 }
333
334 error = check_compiler_feature_info(deviceID, context, test_macro_name,
335 compiler_status);
336 if (error != CL_SUCCESS)
337 {
338 return error;
339 }
340
341 return feature_macro_verify_results(test_macro_name, api_status,
342 compiler_status, supported);
343 }
344
test_feature_macro_device_enqueue(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)345 int test_feature_macro_device_enqueue(cl_device_id deviceID, cl_context context,
346 std::string test_macro_name,
347 cl_bool& supported)
348 {
349 cl_int error = TEST_FAIL;
350 cl_bool api_status;
351 cl_bool compiler_status;
352 log_info("\n%s ...\n", test_macro_name.c_str());
353 error = check_api_feature_info_capabilities<
354 cl_device_device_enqueue_capabilities>(
355 deviceID, context, api_status, CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES,
356 CL_DEVICE_QUEUE_SUPPORTED);
357 if (error != CL_SUCCESS)
358 {
359 return error;
360 }
361
362 error = check_compiler_feature_info(deviceID, context, test_macro_name,
363 compiler_status);
364 if (error != CL_SUCCESS)
365 {
366 return error;
367 }
368
369 return feature_macro_verify_results(test_macro_name, api_status,
370 compiler_status, supported);
371 }
372
test_feature_macro_generic_address_space(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)373 int test_feature_macro_generic_address_space(cl_device_id deviceID,
374 cl_context context,
375 std::string test_macro_name,
376 cl_bool& supported)
377 {
378 cl_int error = TEST_FAIL;
379 cl_bool api_status;
380 cl_bool compiler_status;
381 log_info("\n%s ...\n", test_macro_name.c_str());
382 error = check_api_feature_info_support(
383 deviceID, context, api_status, CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT);
384 if (error != CL_SUCCESS)
385 {
386 return error;
387 }
388
389 error = check_compiler_feature_info(deviceID, context, test_macro_name,
390 compiler_status);
391 if (error != CL_SUCCESS)
392 {
393 return error;
394 }
395
396 return feature_macro_verify_results(test_macro_name, api_status,
397 compiler_status, supported);
398 }
399
test_feature_macro_pipes(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)400 int test_feature_macro_pipes(cl_device_id deviceID, cl_context context,
401 std::string test_macro_name, cl_bool& supported)
402 {
403 cl_int error = TEST_FAIL;
404 cl_bool api_status;
405 cl_bool compiler_status;
406 log_info("\n%s ...\n", test_macro_name.c_str());
407 error = check_api_feature_info_support(deviceID, context, api_status,
408 CL_DEVICE_PIPE_SUPPORT);
409 if (error != CL_SUCCESS)
410 {
411 return error;
412 }
413
414 error = check_compiler_feature_info(deviceID, context, test_macro_name,
415 compiler_status);
416 if (error != CL_SUCCESS)
417 {
418 return error;
419 }
420
421 return feature_macro_verify_results(test_macro_name, api_status,
422 compiler_status, supported);
423 }
424
test_feature_macro_program_scope_global_variables(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)425 int test_feature_macro_program_scope_global_variables(
426 cl_device_id deviceID, cl_context context, std::string test_macro_name,
427 cl_bool& supported)
428 {
429 cl_int error = TEST_FAIL;
430 cl_bool api_status;
431 cl_bool compiler_status;
432 log_info("\n%s ...\n", test_macro_name.c_str());
433 error = check_api_feature_info_number<size_t>(
434 deviceID, context, api_status, CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE);
435 if (error != CL_SUCCESS)
436 {
437 return error;
438 }
439
440 error = check_compiler_feature_info(deviceID, context, test_macro_name,
441 compiler_status);
442 if (error != CL_SUCCESS)
443 {
444 return error;
445 }
446
447 return feature_macro_verify_results(test_macro_name, api_status,
448 compiler_status, supported);
449 }
450
test_feature_macro_read_write_images(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)451 int test_feature_macro_read_write_images(cl_device_id deviceID,
452 cl_context context,
453 std::string test_macro_name,
454 cl_bool& supported)
455 {
456 cl_int error = TEST_FAIL;
457 cl_bool api_status;
458 cl_bool compiler_status;
459 log_info("\n%s ...\n", test_macro_name.c_str());
460 error = check_api_feature_info_number<cl_uint>(
461 deviceID, context, api_status, CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS);
462 if (error != CL_SUCCESS)
463 {
464 return error;
465 }
466
467 error = check_compiler_feature_info(deviceID, context, test_macro_name,
468 compiler_status);
469 if (error != CL_SUCCESS)
470 {
471 return error;
472 }
473
474 return feature_macro_verify_results(test_macro_name, api_status,
475 compiler_status, supported);
476 }
477
test_feature_macro_subgroups(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)478 int test_feature_macro_subgroups(cl_device_id deviceID, cl_context context,
479 std::string test_macro_name,
480 cl_bool& supported)
481 {
482 cl_int error = TEST_FAIL;
483 cl_bool api_status;
484 cl_bool compiler_status;
485 log_info("\n%s ...\n", test_macro_name.c_str());
486 error = check_api_feature_info_number<cl_uint>(
487 deviceID, context, api_status, CL_DEVICE_MAX_NUM_SUB_GROUPS);
488 if (error != CL_SUCCESS)
489 {
490 return error;
491 }
492
493 error = check_compiler_feature_info(deviceID, context, test_macro_name,
494 compiler_status);
495 if (error != CL_SUCCESS)
496 {
497 return error;
498 }
499
500 return feature_macro_verify_results(test_macro_name, api_status,
501 compiler_status, supported);
502 }
503
test_feature_macro_work_group_collective_functions(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)504 int test_feature_macro_work_group_collective_functions(
505 cl_device_id deviceID, cl_context context, std::string test_macro_name,
506 cl_bool& supported)
507 {
508 cl_int error = TEST_FAIL;
509 cl_bool api_status;
510 cl_bool compiler_status;
511 log_info("\n%s ...\n", test_macro_name.c_str());
512 error = check_api_feature_info_support(
513 deviceID, context, api_status,
514 CL_DEVICE_WORK_GROUP_COLLECTIVE_FUNCTIONS_SUPPORT);
515 if (error != CL_SUCCESS)
516 {
517 return error;
518 }
519
520 error = check_compiler_feature_info(deviceID, context, test_macro_name,
521 compiler_status);
522 if (error != CL_SUCCESS)
523 {
524 return error;
525 }
526
527 return feature_macro_verify_results(test_macro_name, api_status,
528 compiler_status, supported);
529 }
530
test_feature_macro_images(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)531 int test_feature_macro_images(cl_device_id deviceID, cl_context context,
532 std::string test_macro_name, cl_bool& supported)
533 {
534 cl_int error = TEST_FAIL;
535 cl_bool api_status;
536 cl_bool compiler_status;
537 log_info("\n%s ...\n", test_macro_name.c_str());
538 error = check_api_feature_info_support(deviceID, context, api_status,
539 CL_DEVICE_IMAGE_SUPPORT);
540 if (error != CL_SUCCESS)
541 {
542 return error;
543 }
544
545 error = check_compiler_feature_info(deviceID, context, test_macro_name,
546 compiler_status);
547 if (error != CL_SUCCESS)
548 {
549 return error;
550 }
551
552 return feature_macro_verify_results(test_macro_name, api_status,
553 compiler_status, supported);
554 }
555
test_feature_macro_fp64(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)556 int test_feature_macro_fp64(cl_device_id deviceID, cl_context context,
557 std::string test_macro_name, cl_bool& supported)
558 {
559 cl_int error = TEST_FAIL;
560 cl_bool api_status;
561 cl_bool compiler_status;
562 log_info("\n%s ...\n", test_macro_name.c_str());
563 error = check_api_feature_info_capabilities<cl_device_fp_config>(
564 deviceID, context, api_status, CL_DEVICE_DOUBLE_FP_CONFIG,
565 CL_FP_FMA | CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN | CL_FP_DENORM);
566 if (error != CL_SUCCESS)
567 {
568 return error;
569 }
570
571 error = check_compiler_feature_info(deviceID, context, test_macro_name,
572 compiler_status);
573 if (error != CL_SUCCESS)
574 {
575 return error;
576 }
577
578 return feature_macro_verify_results(test_macro_name, api_status,
579 compiler_status, supported);
580 }
581
test_feature_macro_integer_dot_product_input_4x8bit_packed(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)582 int test_feature_macro_integer_dot_product_input_4x8bit_packed(
583 cl_device_id deviceID, cl_context context, std::string test_macro_name,
584 cl_bool& supported)
585 {
586 cl_int error = TEST_FAIL;
587 cl_bool api_status;
588 cl_bool compiler_status;
589 log_info("\n%s ...\n", test_macro_name.c_str());
590
591 if (!is_extension_available(deviceID, "cl_khr_integer_dot_product"))
592 {
593 supported = false;
594 return TEST_PASS;
595 }
596
597 error = check_api_feature_info_capabilities<
598 cl_device_integer_dot_product_capabilities_khr>(
599 deviceID, context, api_status,
600 CL_DEVICE_INTEGER_DOT_PRODUCT_CAPABILITIES_KHR,
601 CL_DEVICE_INTEGER_DOT_PRODUCT_INPUT_4x8BIT_PACKED_KHR);
602 if (error != CL_SUCCESS)
603 {
604 return error;
605 }
606
607 error = check_compiler_feature_info(deviceID, context, test_macro_name,
608 compiler_status);
609 if (error != CL_SUCCESS)
610 {
611 return error;
612 }
613
614 return feature_macro_verify_results(test_macro_name, api_status,
615 compiler_status, supported);
616 }
617
test_feature_macro_integer_dot_product_input_4x8bit(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)618 int test_feature_macro_integer_dot_product_input_4x8bit(
619 cl_device_id deviceID, cl_context context, std::string test_macro_name,
620 cl_bool& supported)
621 {
622 cl_int error = TEST_FAIL;
623 cl_bool api_status;
624 cl_bool compiler_status;
625 log_info("\n%s ...\n", test_macro_name.c_str());
626
627 if (!is_extension_available(deviceID, "cl_khr_integer_dot_product"))
628 {
629 supported = false;
630 return TEST_PASS;
631 }
632
633 error = check_api_feature_info_capabilities<
634 cl_device_integer_dot_product_capabilities_khr>(
635 deviceID, context, api_status,
636 CL_DEVICE_INTEGER_DOT_PRODUCT_CAPABILITIES_KHR,
637 CL_DEVICE_INTEGER_DOT_PRODUCT_INPUT_4x8BIT_KHR);
638 if (error != CL_SUCCESS)
639 {
640 return error;
641 }
642
643 error = check_compiler_feature_info(deviceID, context, test_macro_name,
644 compiler_status);
645 if (error != CL_SUCCESS)
646 {
647 return error;
648 }
649
650 return feature_macro_verify_results(test_macro_name, api_status,
651 compiler_status, supported);
652 }
653
test_feature_macro_int64(cl_device_id deviceID,cl_context context,std::string test_macro_name,cl_bool & supported)654 int test_feature_macro_int64(cl_device_id deviceID, cl_context context,
655 std::string test_macro_name, cl_bool& supported)
656 {
657 cl_int error = TEST_FAIL;
658 cl_bool api_status;
659 cl_bool compiler_status;
660 cl_int full_profile = 0;
661 log_info("\n%s ...\n", test_macro_name.c_str());
662 size_t ret_len;
663 char profile[32] = { 0 };
664 error = clGetDeviceInfo(deviceID, CL_DEVICE_PROFILE, sizeof(profile),
665 profile, &ret_len);
666 test_error(error, "clGetDeviceInfo(CL_DEVICE_PROFILE) failed");
667 if (ret_len < sizeof(profile) && strcmp(profile, "FULL_PROFILE") == 0)
668 {
669 full_profile = 1;
670 }
671 else if (ret_len < sizeof(profile)
672 && strcmp(profile, "EMBEDDED_PROFILE") == 0)
673 {
674 full_profile = 0;
675 }
676 else
677 {
678 log_error("Unknown device profile: %s\n", profile);
679 return TEST_FAIL;
680 }
681
682 if (full_profile)
683 {
684 api_status = CL_TRUE;
685 }
686 else
687 {
688 if (is_extension_available(deviceID, "cles_khr_int64"))
689 {
690 api_status = CL_TRUE;
691 }
692 else
693 {
694 cl_bool double_supported = CL_FALSE;
695 error = check_api_feature_info_capabilities<cl_device_fp_config>(
696 deviceID, context, double_supported, CL_DEVICE_DOUBLE_FP_CONFIG,
697 CL_FP_FMA | CL_FP_ROUND_TO_NEAREST | CL_FP_INF_NAN
698 | CL_FP_DENORM);
699 test_error(error, "checking CL_DEVICE_DOUBLE_FP_CONFIG failed");
700 if (double_supported == CL_FALSE)
701 {
702 api_status = CL_FALSE;
703 }
704 else
705 {
706 log_error("FP double type is supported and cles_khr_int64 "
707 "extension not supported\n");
708 return TEST_FAIL;
709 }
710 }
711 }
712
713 error = check_compiler_feature_info(deviceID, context, test_macro_name,
714 compiler_status);
715 if (error != CL_SUCCESS)
716 {
717 return error;
718 }
719
720 return feature_macro_verify_results(test_macro_name, api_status,
721 compiler_status, supported);
722 }
723
test_consistency_c_features_list(cl_device_id deviceID,std::vector<std::string> vec_to_cmp)724 int test_consistency_c_features_list(cl_device_id deviceID,
725 std::vector<std::string> vec_to_cmp)
726 {
727 log_info("\nComparison list of features: CL_DEVICE_OPENCL_C_FEATURES vs "
728 "API/compiler queries.\n");
729 cl_int error;
730 size_t config_size;
731 std::vector<cl_name_version> vec_device_feature;
732 std::vector<std::string> vec_device_feature_names;
733 error = clGetDeviceInfo(deviceID, CL_DEVICE_OPENCL_C_FEATURES, 0, NULL,
734 &config_size);
735
736 test_error(
737 error,
738 "clGetDeviceInfo asking for CL_DEVICE_OPENCL_C_FEATURES failed.\n");
739 if (config_size == 0)
740 {
741 log_info("Empty list of CL_DEVICE_OPENCL_C_FEATURES returned by "
742 "clGetDeviceInfo on this device.\n");
743 }
744 else
745 {
746 int vec_elements = config_size / sizeof(cl_name_version);
747 vec_device_feature.resize(vec_elements);
748 error = clGetDeviceInfo(deviceID, CL_DEVICE_OPENCL_C_FEATURES,
749 config_size, vec_device_feature.data(), 0);
750 test_error(
751 error,
752 "clGetDeviceInfo asking for CL_DEVICE_OPENCL_C_FEATURES failed.\n");
753 }
754 for (auto each_f : vec_device_feature)
755 {
756 vec_device_feature_names.push_back(each_f.name);
757 }
758 sort(vec_to_cmp.begin(), vec_to_cmp.end());
759 sort(vec_device_feature_names.begin(), vec_device_feature_names.end());
760
761 log_info(
762 "Supported features based on CL_DEVICE_OPENCL_C_FEATURES API query:\n");
763 for (auto each_f : vec_device_feature_names)
764 {
765 log_info("%s\n", each_f.c_str());
766 }
767
768 log_info("\nSupported features based on queries to API/compiler :\n");
769
770 for (auto each_f : vec_to_cmp)
771 {
772 log_info("%s\n", each_f.c_str());
773 }
774
775 for (auto each_f : vec_to_cmp)
776 {
777 if (find(vec_device_feature_names.begin(),
778 vec_device_feature_names.end(), each_f)
779 == vec_device_feature_names.end())
780 {
781 log_info("Comparison list of features - failed - missing %s\n",
782 each_f.c_str());
783 return TEST_FAIL;
784 }
785 }
786
787 log_info("Comparison list of features - passed\n");
788
789 return error;
790 }
791
792 #define NEW_FEATURE_MACRO_TEST(feat) \
793 test_macro_name = "__opencl_c_" #feat; \
794 error |= test_feature_macro_##feat(deviceID, context, test_macro_name, \
795 supported); \
796 if (supported) supported_features_vec.push_back(test_macro_name);
797
798
test_features_macro(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)799 int test_features_macro(cl_device_id deviceID, cl_context context,
800 cl_command_queue queue, int num_elements)
801 {
802
803 // Note: Not checking that the feature array is empty for the compiler not
804 // available case because the specification says "For devices that do not
805 // support compilation from OpenCL C source, this query may return an empty
806 // array." It "may" return an empty array implies that an implementation
807 // also "may not".
808 check_compiler_available(deviceID);
809
810 int error = TEST_PASS;
811 cl_bool supported = CL_FALSE;
812 std::string test_macro_name = "";
813 std::vector<std::string> supported_features_vec;
814 NEW_FEATURE_MACRO_TEST(program_scope_global_variables);
815 NEW_FEATURE_MACRO_TEST(3d_image_writes);
816 NEW_FEATURE_MACRO_TEST(atomic_order_acq_rel);
817 NEW_FEATURE_MACRO_TEST(atomic_order_seq_cst);
818 NEW_FEATURE_MACRO_TEST(atomic_scope_device);
819 NEW_FEATURE_MACRO_TEST(atomic_scope_all_devices);
820 NEW_FEATURE_MACRO_TEST(device_enqueue);
821 NEW_FEATURE_MACRO_TEST(generic_address_space);
822 NEW_FEATURE_MACRO_TEST(pipes);
823 NEW_FEATURE_MACRO_TEST(read_write_images);
824 NEW_FEATURE_MACRO_TEST(subgroups);
825 NEW_FEATURE_MACRO_TEST(work_group_collective_functions);
826 NEW_FEATURE_MACRO_TEST(images);
827 NEW_FEATURE_MACRO_TEST(fp64);
828 NEW_FEATURE_MACRO_TEST(int64);
829 NEW_FEATURE_MACRO_TEST(integer_dot_product_input_4x8bit);
830 NEW_FEATURE_MACRO_TEST(integer_dot_product_input_4x8bit_packed);
831
832 error |= test_consistency_c_features_list(deviceID, supported_features_vec);
833
834 return error;
835 }
836