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