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_spirv_linkonce_odr",
80 "cl_khr_semaphore",
81 "cl_khr_external_semaphore",
82 "cl_khr_external_semaphore_win32",
83 "cl_khr_external_semaphore_sync_fd",
84 "cl_khr_external_semaphore_opaque_fd",
85 "cl_khr_external_memory",
86 "cl_khr_external_memory_win32",
87 "cl_khr_external_memory_opaque_fd",
88 "cl_khr_command_buffer",
89 "cl_khr_command_buffer_mutable_dispatch",
90 };
91
92 size_t num_known_extensions = ARRAY_SIZE(known_extensions);
93 size_t first_API_extension = 31;
94
95 const char *known_embedded_extensions[] = {
96 "cles_khr_int64",
97 NULL
98 };
99
100 typedef enum
101 {
102 kUnsupported_extension = -1,
103 kVendor_extension = 0,
104 kLanguage_extension = 1,
105 kAPI_extension = 2
106 }Extension_Type;
107
108 const char *kernel_strings[] = {
109 "kernel void test(global int *defines)\n{\n",
110 "#pragma OPENCL EXTENSION %s : enable\n",
111 "#ifdef %s\n"
112 " defines[%d] = 1;\n"
113 "#else\n"
114 " defines[%d] = 0;\n"
115 "#endif\n",
116 "#pragma OPENCL EXTENSION %s : disable\n\n",
117 "}\n"
118 };
119
test_compiler_defines_for_extensions(cl_device_id device,cl_context context,cl_command_queue queue,int n_elems)120 int test_compiler_defines_for_extensions(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems )
121 {
122
123 int error;
124 int total_errors = 0;
125
126
127 // Get the extensions string for the device
128 size_t size;
129 error = clGetDeviceInfo(device, CL_DEVICE_EXTENSIONS, 0, NULL, &size);
130 test_error(error, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS size failed");
131
132 char *extensions = (char*)malloc(sizeof(char)*(size + 1));
133 if (extensions == 0) {
134 log_error("Failed to allocate memory for extensions string.\n");
135 return -1;
136 }
137 memset( extensions, CHAR_MIN, sizeof(char)*(size+1) );
138
139 error = clGetDeviceInfo(device, CL_DEVICE_EXTENSIONS, sizeof(char)*size, extensions, NULL);
140 test_error(error, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS failed");
141
142 // Check to make sure the extension string is NUL terminated.
143 if( extensions[size] != CHAR_MIN )
144 {
145 test_error( -1, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS wrote past the end of the array!" );
146 return -1;
147 }
148 extensions[size] = '\0'; // set last char to NUL to avoid problems with string functions later
149
150 // test for termination with '\0'
151 size_t stringSize = strlen( extensions );
152 if( stringSize == size )
153 {
154 test_error( -1, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS is not NUL terminated!" );
155 return -1;
156 }
157
158 // Break up the extensions
159 log_info("Device reports the following extensions:\n");
160 char *extensions_supported[1024];
161 Extension_Type extension_type[1024];
162 int num_of_supported_extensions = 0;
163 char *currentP = extensions;
164
165 memset( extension_type, 0, sizeof( extension_type) );
166
167 // loop over extension string
168 while (currentP != extensions + stringSize)
169 {
170 // skip leading white space
171 while( *currentP == ' ' )
172 currentP++;
173
174 // Exit if end of string
175 if( *currentP == '\0' )
176 {
177 if( currentP != extensions + stringSize)
178 {
179 test_error( -1, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS contains a NUL in the middle of the string!" );
180 return -1;
181 }
182 break;
183 }
184
185 // Not space, not end of string, so extension
186 char *start = currentP; // start of extension name
187
188 // loop looking for the end
189 while (*currentP != ' ' && currentP != extensions + stringSize)
190 {
191 // check for non-space white space in the extension name
192 if( isspace(*currentP) )
193 {
194 test_error( -1, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS contains a non-space whitespace in an extension name!" );
195 return -1;
196 }
197 currentP++;
198 }
199
200 // record the extension name
201 uintptr_t extension_length = (uintptr_t) currentP - (uintptr_t) start;
202 extensions_supported[ num_of_supported_extensions ] = (char*) malloc( (extension_length + 1) * sizeof( char ) );
203 if( NULL == extensions_supported[ num_of_supported_extensions ] )
204 {
205 log_error( "Error: unable to allocate memory to hold extension name: %ld chars\n", extension_length );
206 return -1;
207 }
208 memcpy( extensions_supported[ num_of_supported_extensions ], start, extension_length * sizeof( char ) );
209 extensions_supported[ num_of_supported_extensions ][extension_length] = '\0';
210
211 // If the extension is a cl_khr extension, make sure it is an approved cl_khr extension -- looking for misspellings here
212 if( extensions_supported[ num_of_supported_extensions ][0] == 'c' &&
213 extensions_supported[ num_of_supported_extensions ][1] == 'l' &&
214 extensions_supported[ num_of_supported_extensions ][2] == '_' &&
215 extensions_supported[ num_of_supported_extensions ][3] == 'k' &&
216 extensions_supported[ num_of_supported_extensions ][4] == 'h' &&
217 extensions_supported[ num_of_supported_extensions ][5] == 'r' &&
218 extensions_supported[ num_of_supported_extensions ][6] == '_' )
219 {
220 size_t ii;
221 for( ii = 0; ii < num_known_extensions; ii++ )
222 {
223 if( 0 == strcmp( known_extensions[ii], extensions_supported[ num_of_supported_extensions ] ) )
224 break;
225 }
226 if( ii == num_known_extensions )
227 {
228 log_error( "FAIL: Extension %s is not in the list of approved Khronos extensions!", extensions_supported[ num_of_supported_extensions ] );
229 return -1;
230 }
231 }
232 // Is it an embedded extension?
233 else if( memcmp( extensions_supported[ num_of_supported_extensions ], "cles_khr_", 9 ) == 0 )
234 {
235 // Yes, but is it a known one?
236 size_t ii;
237 for( ii = 0; known_embedded_extensions[ ii ] != NULL; ii++ )
238 {
239 if( strcmp( known_embedded_extensions[ ii ], extensions_supported[ num_of_supported_extensions ] ) == 0 )
240 break;
241 }
242 if( known_embedded_extensions[ ii ] == NULL )
243 {
244 log_error( "FAIL: Extension %s is not in the list of approved Khronos embedded extensions!", extensions_supported[ num_of_supported_extensions ] );
245 return -1;
246 }
247
248 // It's approved, but are we even an embedded system?
249 char profileStr[128] = "";
250 error = clGetDeviceInfo( device, CL_DEVICE_PROFILE, sizeof( profileStr ), &profileStr, NULL );
251 test_error( error, "Unable to get CL_DEVICE_PROFILE to validate embedded extension name" );
252
253 if( strcmp( profileStr, "EMBEDDED_PROFILE" ) != 0 )
254 {
255 log_error( "FAIL: Extension %s is an approved embedded extension, but on a non-embedded profile!", extensions_supported[ num_of_supported_extensions ] );
256 return -1;
257 }
258 }
259 else
260 { // All other extensions must be of the form cl_<vendor_name>_<name>
261 if( extensions_supported[ num_of_supported_extensions ][0] != 'c' ||
262 extensions_supported[ num_of_supported_extensions ][1] != 'l' ||
263 extensions_supported[ num_of_supported_extensions ][2] != '_' )
264 {
265 log_error( "FAIL: Extension %s doesn't start with \"cl_\"!", extensions_supported[ num_of_supported_extensions ] );
266 return -1;
267 }
268
269 if( extensions_supported[ num_of_supported_extensions ][3] == '_' || extensions_supported[ num_of_supported_extensions ][3] == '\0' )
270 {
271 log_error( "FAIL: Vendor name is missing in extension %s!", extensions_supported[ num_of_supported_extensions ] );
272 return -1;
273 }
274
275 // look for the second underscore for name
276 char *p = extensions_supported[ num_of_supported_extensions ] + 4;
277 while( *p != '\0' && *p != '_' )
278 p++;
279
280 if( *p != '_' || p[1] == '\0')
281 {
282 log_error( "FAIL: extension name is missing in extension %s!", extensions_supported[ num_of_supported_extensions ] );
283 return -1;
284 }
285 }
286
287
288 num_of_supported_extensions++;
289 }
290
291 // Build a list of the known extensions that are not supported by the device
292 char *extensions_not_supported[1024];
293 int num_not_supported_extensions = 0;
294 for( int i = 0; i < num_of_supported_extensions; i++ )
295 {
296 int is_supported = 0;
297 for( size_t j = 0; j < num_known_extensions; j++ )
298 {
299 if( strcmp( extensions_supported[ i ], known_extensions[ j ] ) == 0 )
300 {
301 extension_type[ i ] = ( j < first_API_extension ) ? kLanguage_extension : kAPI_extension;
302 is_supported = 1;
303 break;
304 }
305 }
306 if( !is_supported )
307 {
308 for( int j = 0; known_embedded_extensions[ j ] != NULL; j++ )
309 {
310 if( strcmp( extensions_supported[ i ], known_embedded_extensions[ j ] ) == 0 )
311 {
312 extension_type[ i ] = kLanguage_extension;
313 is_supported = 1;
314 break;
315 }
316 }
317 }
318 if (!is_supported) {
319 extensions_not_supported[num_not_supported_extensions] = (char*)malloc(strlen(extensions_supported[i])+1);
320 strcpy(extensions_not_supported[num_not_supported_extensions], extensions_supported[i]);
321 num_not_supported_extensions++;
322 }
323 }
324
325 for (int i=0; i<num_of_supported_extensions; i++) {
326 log_info("%40s -- Supported\n", extensions_supported[i]);
327 }
328 for (int i=0; i<num_not_supported_extensions; i++) {
329 log_info("%40s -- Not Supported\n", extensions_not_supported[i]);
330 }
331
332 // Build the kernel
333 char *kernel_code = (char *)malloc(
334 1
335 + 1025 * 256
336 * (num_not_supported_extensions + num_of_supported_extensions));
337 memset(
338 kernel_code, 0,
339 1
340 + 1025 * 256
341 * (num_not_supported_extensions + num_of_supported_extensions));
342
343 int i, index = 0;
344 strcat(kernel_code, kernel_strings[0]);
345 for (i=0; i<num_of_supported_extensions; i++, index++) {
346
347 if (extension_type[i] == kLanguage_extension)
348 sprintf(kernel_code + strlen(kernel_code), kernel_strings[1], extensions_supported[i]);
349
350 sprintf(kernel_code + strlen(kernel_code), kernel_strings[2], extensions_supported[i], index, index );
351
352 if (extension_type[i] == kLanguage_extension)
353 sprintf(kernel_code + strlen(kernel_code), kernel_strings[3], extensions_supported[i] );
354 }
355 for ( i = 0; i<num_not_supported_extensions; i++, index++) {
356 sprintf(kernel_code + strlen(kernel_code), kernel_strings[2], extensions_not_supported[i], index, index );
357 }
358 strcat(kernel_code, kernel_strings[4]);
359
360 // Now we need to execute the kernel
361 clMemWrapper defines;
362 cl_int *data;
363 clProgramWrapper program;
364 clKernelWrapper kernel;
365
366 error = create_single_kernel_helper(context, &program, &kernel, 1,
367 (const char **)&kernel_code, "test");
368 test_error(error, "create_single_kernel_helper failed");
369
370 data = (cl_int*)malloc(sizeof(cl_int)*(num_not_supported_extensions+num_of_supported_extensions));
371 memset(data, 0, sizeof(cl_int)*(num_not_supported_extensions+num_of_supported_extensions));
372 defines = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
373 sizeof(cl_int)*(num_not_supported_extensions+num_of_supported_extensions), data, &error);
374 test_error(error, "clCreateBuffer failed");
375
376 error = clSetKernelArg(kernel, 0, sizeof(defines), &defines);
377 test_error(error, "clSetKernelArg failed");
378
379 size_t global_size = 1;
380 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, NULL, 0, NULL, NULL);
381 test_error(error, "clEnqueueNDRangeKernel failed");
382
383 error = clEnqueueReadBuffer(queue, defines, CL_TRUE, 0, sizeof(cl_int)*(num_not_supported_extensions+num_of_supported_extensions),
384 data, 0, NULL, NULL);
385 test_error(error, "clEnqueueReadBuffer failed");
386
387 // Report what the compiler reported
388 log_info("\nCompiler reported the following extensions defined in the OpenCL C kernel environment:\n");
389 index = 0;
390 int total_supported = 0;
391 for (int i=0; i<num_of_supported_extensions; i++, index++) {
392 if (data[index] == 1) {
393 log_info("\t%s\n", extensions_supported[i]);
394 total_supported++;
395 }
396 }
397 for (int i=0; i<num_not_supported_extensions; i++, index++) {
398 if (data[index] == 1) {
399 log_info("\t%s\n", extensions_not_supported[i]);
400 total_supported++;
401 }
402 }
403 if (total_supported == 0)
404 log_info("\t(none)\n");
405
406 // Count the errors
407 index = 0;
408 int unknown = 0;
409 for ( i=0; i<num_of_supported_extensions; i++)
410 {
411 if (data[i] != 1)
412 {
413 switch( extension_type[i] )
414 {
415 case kLanguage_extension:
416 log_error("ERROR: Supported extension %s not defined in kernel.\n", extensions_supported[i]);
417 total_errors++;
418 break;
419 case kVendor_extension:
420 unknown++;
421 break;
422 case kAPI_extension:
423 break;
424 default:
425 log_error( "ERROR: internal test error in extension detection. This is probably a bug in the test.\n" );
426 break;
427 }
428 }
429 }
430
431 if(unknown)
432 {
433 log_info( "\nThe following non-KHR extensions are supported but do not add a preprocessor symbol to OpenCL C.\n" );
434 for (int z=0; z<num_of_supported_extensions; z++)
435 {
436 if (data[z] != 1 && extension_type[z] == kVendor_extension )
437 log_info( "\t%s\n", extensions_supported[z]);
438 }
439 }
440
441 for ( ; i<num_not_supported_extensions; i++) {
442 if (data[i] != 0) {
443 log_error("ERROR: Unsupported extension %s is defined in kernel.\n", extensions_not_supported[i]);
444 total_errors++;
445 }
446 }
447 log_info("\n");
448
449 // cleanup
450 free(data);
451 free(kernel_code);
452 for(i=0; i<num_of_supported_extensions; i++) {
453 free(extensions_supported[i]);
454 }
455 free(extensions);
456
457 if (total_errors)
458 return -1;
459 return 0;
460 }
461