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 "testBase.h"
17 #include <limits.h>
18 #include <ctype.h>
19 #ifndef _WIN32
20 #include <unistd.h>
21 #endif
22
23 // List should follow order in the extension spec
24 const char *known_extensions[] = {
25 "cl_khr_byte_addressable_store",
26 "cl_khr_3d_image_writes",
27 "cl_khr_fp16",
28 "cl_khr_fp64",
29 "cl_khr_global_int32_base_atomics",
30 "cl_khr_global_int32_extended_atomics",
31 "cl_khr_local_int32_base_atomics",
32 "cl_khr_local_int32_extended_atomics",
33 "cl_khr_int64_base_atomics",
34 "cl_khr_int64_extended_atomics",
35 "cl_khr_select_fprounding_mode",
36 "cl_khr_depth_images",
37 "cl_khr_gl_depth_images",
38 "cl_khr_gl_msaa_sharing",
39 "cl_khr_device_enqueue_local_arg_types",
40 "cl_khr_subgroups",
41 "cl_khr_mipmap_image",
42 "cl_khr_mipmap_image_writes",
43 "cl_khr_srgb_image_writes",
44 "cl_khr_subgroup_named_barrier",
45 "cl_khr_extended_async_copies",
46 "cl_khr_subgroup_extended_types",
47 "cl_khr_subgroup_non_uniform_vote",
48 "cl_khr_subgroup_ballot",
49 "cl_khr_subgroup_non_uniform_arithmetic",
50 "cl_khr_subgroup_shuffle",
51 "cl_khr_subgroup_shuffle_relative",
52 "cl_khr_subgroup_clustered_reduce",
53 "cl_khr_extended_bit_ops",
54 "cl_khr_integer_dot_product",
55 "cl_khr_subgroup_rotate",
56 // API-only extensions after this point. If you add above here, modify
57 // first_API_extension below.
58 "cl_khr_icd",
59 "cl_khr_gl_sharing",
60 "cl_khr_gl_event",
61 "cl_khr_d3d10_sharing",
62 "cl_khr_d3d11_sharing",
63 "cl_khr_dx9_media_sharing",
64 "cl_khr_egl_event",
65 "cl_khr_egl_image",
66 "cl_khr_image2d_from_buffer",
67 "cl_khr_spir",
68 "cl_khr_il_program",
69 "cl_khr_create_command_queue",
70 "cl_khr_initialize_memory",
71 "cl_khr_terminate_context",
72 "cl_khr_priority_hints",
73 "cl_khr_throttle_hints",
74 "cl_khr_spirv_no_integer_wrap_decoration",
75 "cl_khr_extended_versioning",
76 "cl_khr_device_uuid",
77 "cl_khr_pci_bus_info",
78 "cl_khr_suggested_local_work_size",
79 "cl_khr_expect_assume",
80 "cl_khr_spirv_linkonce_odr",
81 "cl_khr_semaphore",
82 "cl_khr_external_semaphore",
83 "cl_khr_external_semaphore_win32",
84 "cl_khr_external_semaphore_sync_fd",
85 "cl_khr_external_semaphore_opaque_fd",
86 "cl_khr_external_semaphore_dx_fence",
87 "cl_khr_external_memory",
88 "cl_khr_external_memory_win32",
89 "cl_khr_external_memory_opaque_fd",
90 "cl_khr_external_memory_dx",
91 "cl_khr_external_memory_dma_buf",
92 "cl_khr_command_buffer",
93 "cl_khr_command_buffer_mutable_dispatch",
94 };
95
96 size_t num_known_extensions = ARRAY_SIZE(known_extensions);
97 size_t first_API_extension = 31;
98
99 const char *known_embedded_extensions[] = {
100 "cles_khr_int64",
101 NULL
102 };
103
104 typedef enum
105 {
106 kUnsupported_extension = -1,
107 kVendor_extension = 0,
108 kLanguage_extension = 1,
109 kAPI_extension = 2
110 }Extension_Type;
111
112 const char *kernel_strings[] = {
113 "kernel void test(global int *defines)\n{\n",
114 "#pragma OPENCL EXTENSION %s : enable\n",
115 "#ifdef %s\n"
116 " defines[%d] = 1;\n"
117 "#else\n"
118 " defines[%d] = 0;\n"
119 "#endif\n",
120 "#pragma OPENCL EXTENSION %s : disable\n\n",
121 "}\n"
122 };
123
string_has_prefix(const char * str,const char * prefix)124 bool string_has_prefix(const char *str, const char *prefix)
125 {
126 return strncmp(str, prefix, strlen(prefix)) == 0;
127 }
128
test_compiler_defines_for_extensions(cl_device_id device,cl_context context,cl_command_queue queue,int n_elems)129 int test_compiler_defines_for_extensions(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
130 {
131
132 int error;
133 int total_errors = 0;
134
135
136 // Get the extensions string for the device
137 size_t size;
138 error = clGetDeviceInfo(device, CL_DEVICE_EXTENSIONS, 0, NULL, &size);
139 test_error(error, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS size failed");
140
141 char *extensions = (char*)malloc(sizeof(char)*(size + 1));
142 if (extensions == 0) {
143 log_error("Failed to allocate memory for extensions string.\n");
144 return -1;
145 }
146 memset( extensions, CHAR_MIN, sizeof(char)*(size+1) );
147
148 error = clGetDeviceInfo(device, CL_DEVICE_EXTENSIONS, sizeof(char)*size, extensions, NULL);
149 test_error(error, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS failed");
150
151 // Check to make sure the extension string is NUL terminated.
152 if( extensions[size] != CHAR_MIN )
153 {
154 test_error( -1, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS wrote past the end of the array!" );
155 return -1;
156 }
157 extensions[size] = '\0'; // set last char to NUL to avoid problems with string functions later
158
159 // test for termination with '\0'
160 size_t stringSize = strlen( extensions );
161 if( stringSize == size )
162 {
163 test_error( -1, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS is not NUL terminated!" );
164 return -1;
165 }
166
167 // Break up the extensions
168 log_info("Device reports the following extensions:\n");
169 char *extensions_supported[1024];
170 Extension_Type extension_type[1024];
171 int num_of_supported_extensions = 0;
172 char *currentP = extensions;
173
174 memset( extension_type, 0, sizeof( extension_type) );
175
176 bool failed = false;
177 // loop over extension string
178 while (currentP != extensions + stringSize)
179 {
180 // skip leading white space
181 while( *currentP == ' ' )
182 currentP++;
183
184 // Exit if end of string
185 if( *currentP == '\0' )
186 {
187 if( currentP != extensions + stringSize)
188 {
189 test_error( -1, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS contains a NUL in the middle of the string!" );
190 return -1;
191 }
192 break;
193 }
194
195 // Not space, not end of string, so extension
196 char *start = currentP; // start of extension name
197
198 // loop looking for the end
199 while (*currentP != ' ' && currentP != extensions + stringSize)
200 {
201 // check for non-space white space in the extension name
202 if( isspace(*currentP) )
203 {
204 test_error( -1, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS contains a non-space whitespace in an extension name!" );
205 return -1;
206 }
207 currentP++;
208 }
209
210 // record the extension name
211 uintptr_t extension_length = (uintptr_t) currentP - (uintptr_t) start;
212 char *extension = (char *)malloc((extension_length + 1) * sizeof(char));
213 if (extension == NULL)
214 {
215 log_error( "Error: unable to allocate memory to hold extension name: %ld chars\n", extension_length );
216 return -1;
217 }
218 extensions_supported[num_of_supported_extensions] = extension;
219 memcpy(extension, start, extension_length * sizeof(char));
220 extension[extension_length] = '\0';
221
222 // If the extension is a cl_khr extension, make sure it is an approved cl_khr extension -- looking for misspellings here
223 if (string_has_prefix(extension, "cl_khr_"))
224 {
225 size_t ii;
226 for( ii = 0; ii < num_known_extensions; ii++ )
227 {
228 if (strcmp(known_extensions[ii], extension) == 0) break;
229 }
230 if( ii == num_known_extensions )
231 {
232 log_error("FAIL: Extension %s is not in the list of approved "
233 "Khronos extensions!\n",
234 extension);
235 failed = true;
236 }
237 }
238 // Is it an embedded extension?
239 else if (string_has_prefix(extension, "cles_khr_"))
240 {
241 // Yes, but is it a known one?
242 size_t ii;
243 for( ii = 0; known_embedded_extensions[ ii ] != NULL; ii++ )
244 {
245 if (strcmp(known_embedded_extensions[ii], extension) == 0)
246 break;
247 }
248 if( known_embedded_extensions[ ii ] == NULL )
249 {
250 log_error("FAIL: Extension %s is not in the list of approved "
251 "Khronos embedded extensions!\n",
252 extension);
253 failed = true;
254 }
255 else
256 {
257 // It's approved, but are we even an embedded system?
258 char profileStr[128] = "";
259 error = clGetDeviceInfo(device, CL_DEVICE_PROFILE,
260 sizeof(profileStr), &profileStr, NULL);
261 test_error(error,
262 "Unable to get CL_DEVICE_PROFILE to validate "
263 "embedded extension name");
264
265 if (strcmp(profileStr, "EMBEDDED_PROFILE") != 0)
266 {
267 log_error(
268 "FAIL: Extension %s is an approved embedded extension, "
269 "but on a non-embedded profile!\n",
270 extension);
271 failed = true;
272 }
273 }
274 }
275 else
276 { // All other extensions must be of the form cl_<vendor_name>_<name>
277 if (!string_has_prefix(extension, "cl_"))
278 {
279 log_error("FAIL: Extension %s doesn't start with \"cl_\"!\n",
280 extension);
281 failed = true;
282 }
283 else if (extension[3] == '_' || extension[3] == '\0')
284 {
285 log_error("FAIL: Vendor name is missing in extension %s!\n",
286 extension);
287 failed = true;
288 }
289 else
290 {
291 // look for the second underscore for name
292 char *p = extension + 4;
293 while (*p != '\0' && *p != '_') p++;
294
295 if (*p != '_' || p[1] == '\0')
296 {
297 log_error(
298 "FAIL: extension name is missing in extension %s!\n",
299 extension);
300 failed = true;
301 }
302 }
303 }
304
305
306 num_of_supported_extensions++;
307 }
308
309 if (failed)
310 {
311 return -1;
312 }
313
314 // Build a list of the known extensions that are not supported by the device
315 char *extensions_not_supported[1024];
316 int num_not_supported_extensions = 0;
317 for( int i = 0; i < num_of_supported_extensions; i++ )
318 {
319 int is_supported = 0;
320 for( size_t j = 0; j < num_known_extensions; j++ )
321 {
322 if( strcmp( extensions_supported[ i ], known_extensions[ j ] ) == 0 )
323 {
324 extension_type[ i ] = ( j < first_API_extension ) ? kLanguage_extension : kAPI_extension;
325 is_supported = 1;
326 break;
327 }
328 }
329 if( !is_supported )
330 {
331 for( int j = 0; known_embedded_extensions[ j ] != NULL; j++ )
332 {
333 if( strcmp( extensions_supported[ i ], known_embedded_extensions[ j ] ) == 0 )
334 {
335 extension_type[ i ] = kLanguage_extension;
336 is_supported = 1;
337 break;
338 }
339 }
340 }
341 if (!is_supported) {
342 extensions_not_supported[num_not_supported_extensions] = (char*)malloc(strlen(extensions_supported[i])+1);
343 strcpy(extensions_not_supported[num_not_supported_extensions], extensions_supported[i]);
344 num_not_supported_extensions++;
345 }
346 }
347
348 for (int i=0; i<num_of_supported_extensions; i++) {
349 log_info("%40s -- Supported\n", extensions_supported[i]);
350 }
351 for (int i=0; i<num_not_supported_extensions; i++) {
352 log_info("%40s -- Not Supported\n", extensions_not_supported[i]);
353 }
354
355 // Build the kernel
356 char *kernel_code = (char *)malloc(
357 1
358 + 1025 * 256
359 * (num_not_supported_extensions + num_of_supported_extensions));
360 memset(
361 kernel_code, 0,
362 1
363 + 1025 * 256
364 * (num_not_supported_extensions + num_of_supported_extensions));
365
366 int i, index = 0;
367 strcat(kernel_code, kernel_strings[0]);
368 for (i=0; i<num_of_supported_extensions; i++, index++) {
369
370 if (extension_type[i] == kLanguage_extension)
371 sprintf(kernel_code + strlen(kernel_code), kernel_strings[1], extensions_supported[i]);
372
373 sprintf(kernel_code + strlen(kernel_code), kernel_strings[2], extensions_supported[i], index, index );
374
375 if (extension_type[i] == kLanguage_extension)
376 sprintf(kernel_code + strlen(kernel_code), kernel_strings[3], extensions_supported[i] );
377 }
378 for ( i = 0; i<num_not_supported_extensions; i++, index++) {
379 sprintf(kernel_code + strlen(kernel_code), kernel_strings[2], extensions_not_supported[i], index, index );
380 }
381 strcat(kernel_code, kernel_strings[4]);
382
383 // Now we need to execute the kernel
384 clMemWrapper defines;
385 cl_int *data;
386 clProgramWrapper program;
387 clKernelWrapper kernel;
388
389 error = create_single_kernel_helper(context, &program, &kernel, 1,
390 (const char **)&kernel_code, "test");
391 test_error(error, "create_single_kernel_helper failed");
392
393 data = (cl_int*)malloc(sizeof(cl_int)*(num_not_supported_extensions+num_of_supported_extensions));
394 memset(data, 0, sizeof(cl_int)*(num_not_supported_extensions+num_of_supported_extensions));
395 defines = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
396 sizeof(cl_int)*(num_not_supported_extensions+num_of_supported_extensions), data, &error);
397 test_error(error, "clCreateBuffer failed");
398
399 error = clSetKernelArg(kernel, 0, sizeof(defines), &defines);
400 test_error(error, "clSetKernelArg failed");
401
402 size_t global_size = 1;
403 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, NULL, 0, NULL, NULL);
404 test_error(error, "clEnqueueNDRangeKernel failed");
405
406 error = clEnqueueReadBuffer(queue, defines, CL_TRUE, 0, sizeof(cl_int)*(num_not_supported_extensions+num_of_supported_extensions),
407 data, 0, NULL, NULL);
408 test_error(error, "clEnqueueReadBuffer failed");
409
410 // Report what the compiler reported
411 log_info("\nCompiler reported the following extensions defined in the OpenCL C kernel environment:\n");
412 index = 0;
413 int total_supported = 0;
414 for (int i=0; i<num_of_supported_extensions; i++, index++) {
415 if (data[index] == 1) {
416 log_info("\t%s\n", extensions_supported[i]);
417 total_supported++;
418 }
419 }
420 for (int i=0; i<num_not_supported_extensions; i++, index++) {
421 if (data[index] == 1) {
422 log_info("\t%s\n", extensions_not_supported[i]);
423 total_supported++;
424 }
425 }
426 if (total_supported == 0)
427 log_info("\t(none)\n");
428
429 // Count the errors
430 index = 0;
431 int unknown = 0;
432 for ( i=0; i<num_of_supported_extensions; i++)
433 {
434 if (data[i] != 1)
435 {
436 switch( extension_type[i] )
437 {
438 case kLanguage_extension:
439 log_error("ERROR: Supported extension %s not defined in kernel.\n", extensions_supported[i]);
440 total_errors++;
441 break;
442 case kVendor_extension:
443 unknown++;
444 break;
445 case kAPI_extension:
446 break;
447 default:
448 log_error( "ERROR: internal test error in extension detection. This is probably a bug in the test.\n" );
449 break;
450 }
451 }
452 }
453
454 if(unknown)
455 {
456 log_info( "\nThe following non-KHR extensions are supported but do not add a preprocessor symbol to OpenCL C.\n" );
457 for (int z=0; z<num_of_supported_extensions; z++)
458 {
459 if (data[z] != 1 && extension_type[z] == kVendor_extension )
460 log_info( "\t%s\n", extensions_supported[z]);
461 }
462 }
463
464 for ( ; i<num_not_supported_extensions; i++) {
465 if (data[i] != 0) {
466 log_error("ERROR: Unsupported extension %s is defined in kernel.\n", extensions_not_supported[i]);
467 total_errors++;
468 }
469 }
470 log_info("\n");
471
472 // cleanup
473 free(data);
474 free(kernel_code);
475 for(i=0; i<num_of_supported_extensions; i++) {
476 free(extensions_supported[i]);
477 }
478 free(extensions);
479
480 if (total_errors)
481 return -1;
482 return 0;
483 }
484