• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
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