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 "procs.h"
17 #include <ctype.h>
18 
19 // Test __FILE__, __LINE__, __OPENCL_VERSION__, __OPENCL_C_VERSION__, __ENDIAN_LITTLE__, __ROUNDING_MODE__, __IMAGE_SUPPORT__, __FAST_RELAXED_MATH__
20 // __kernel_exec
21 
22 const char *preprocessor_test = {
23     "#line 2 \"%s\"\n"
24     "__kernel void test( __global int *results, __global char *outFileString, __global char *outRoundingString )\n"
25     "{\n"
26 
27     // Integer preprocessor macros
28     "#ifdef __IMAGE_SUPPORT__\n"
29     "    results[0] =    __IMAGE_SUPPORT__;\n"
30     "#else\n"
31     "    results[0] = 0xf00baa;\n"
32     "#endif\n"
33 
34     "#ifdef __ENDIAN_LITTLE__\n"
35     "    results[1] =    __ENDIAN_LITTLE__;\n"
36     "#else\n"
37     "    results[1] = 0xf00baa;\n"
38     "#endif\n"
39 
40     "#ifdef __OPENCL_VERSION__\n"
41     "    results[2] =    __OPENCL_VERSION__;\n"
42     "#else\n"
43     "    results[2] = 0xf00baa;\n"
44     "#endif\n"
45 
46     "#ifdef __OPENCL_C_VERSION__\n"
47     "    results[3] =    __OPENCL_C_VERSION__;\n"
48     "#else\n"
49     "    results[3] = 0xf00baa;\n"
50     "#endif\n"
51 
52     "#ifdef __LINE__\n"
53     "    results[4] =    __LINE__;\n"
54     "#else\n"
55     "    results[4] = 0xf00baa;\n"
56     "#endif\n"
57 
58 #if 0 // Removed by Affie's request 2/24
59     "#ifdef __FAST_RELAXED_MATH__\n"
60     "    results[5] =    __FAST_RELAXED_MATH__;\n"
61     "#else\n"
62     "    results[5] = 0xf00baa;\n"
63     "#endif\n"
64 #endif
65 
66     "#ifdef __kernel_exec\n"
67     "    results[6] = 1;\n"    // By spec, we can only really evaluate that it is defined, not what it expands to
68     "#else\n"
69     "    results[6] = 0xf00baa;\n"
70     "#endif\n"
71 
72     // String preprocessor macros. Technically, there are strings in OpenCL, but not really.
73     "#ifdef __FILE__\n"
74     "    int i;\n"
75     "    constant char *f = \"\" __FILE__;\n"
76     "   for( i = 0; f[ i ] != 0 && i < 512; i++ )\n"
77     "        outFileString[ i ] = f[ i ];\n"
78     "    outFileString[ i ] = 0;\n"
79     "#else\n"
80     "    outFileString[ 0 ] = 0;\n"
81     "#endif\n"
82 
83     "}\n"
84     };
85 
test_kernel_preprocessor_macros(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)86 int test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
87 {
88     clProgramWrapper program;
89     clKernelWrapper kernel;
90     clMemWrapper streams[ 3 ];
91 
92     int error;
93     size_t    threads[] = {1,1,1};
94 
95     cl_int results[ 7 ];
96     cl_char fileString[ 512 ] = "", roundingString[ 128 ] = "";
97     char programSource[4096];
98     char curFileName[512];
99     char *programPtr = programSource;
100     int i = 0;
101     snprintf(curFileName, 512, "%s", __FILE__);
102 #ifdef _WIN32
103     // Replace "\" with "\\"
104     while(curFileName[i] != '\0') {
105         if (curFileName[i] == '\\') {
106             int j = i + 1;
107             char prev = '\\';
108             while (curFileName[j - 1] != '\0') {
109                 char tmp = curFileName[j];
110                 curFileName[j] = prev;
111                 prev = tmp;
112                 j++;
113             }
114             i++;
115         }
116         i++;
117     }
118 #endif
119     sprintf(programSource,preprocessor_test,curFileName);
120 
121     // Create the kernel
122     if( create_single_kernel_helper( context, &program, &kernel, 1,  (const char **)&programPtr, "test" ) != 0 )
123     {
124         return -1;
125     }
126 
127     /* Create some I/O streams */
128     streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(results),
129                                 NULL, &error);
130     test_error( error, "Creating test array failed" );
131     streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(fileString),
132                                 NULL, &error);
133     test_error( error, "Creating test array failed" );
134     streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
135                                 sizeof(roundingString), NULL, &error);
136     test_error( error, "Creating test array failed" );
137 
138     // Set up and run
139     for( int i = 0; i < 3; i++ )
140     {
141         error = clSetKernelArg( kernel, i, sizeof( streams[i] ), &streams[i] );
142         test_error( error, "Unable to set indexed kernel arguments" );
143     }
144 
145     error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
146     test_error( error, "Kernel execution failed" );
147 
148     error = clEnqueueReadBuffer( queue, streams[0], CL_TRUE, 0, sizeof(results), results, 0, NULL, NULL );
149     test_error( error, "Unable to get result data" );
150     error = clEnqueueReadBuffer( queue, streams[1], CL_TRUE, 0, sizeof(fileString), fileString, 0, NULL, NULL );
151     test_error( error, "Unable to get result data" );
152     error = clEnqueueReadBuffer( queue, streams[2], CL_TRUE, 0, sizeof(roundingString), roundingString, 0, NULL, NULL );
153     test_error( error, "Unable to get result data" );
154 
155 
156     /////// Check the integer results
157 
158     // We need to check these values against what we know is supported on the device
159     if( checkForImageSupport( deviceID ) == 0 )
160     {
161         // If images are supported, the constant should have been defined to the value 1
162         if( results[ 0 ] == 0xf00baa )
163         {
164             log_error( "ERROR: __IMAGE_SUPPORT__ undefined even though images are supported\n" );
165             return -1;
166         }
167         else if( results[ 0 ] != 1 )
168         {
169             log_error( "ERROR: __IMAGE_SUPPORT__ defined, but to the wrong value (defined as %d, spec states it should be 1)\n", (int)results[ 0 ] );
170             return -1;
171         }
172     }
173     else
174     {
175         // If images aren't supported, the constant should be undefined
176         if( results[ 0 ] != 0xf00baa )
177         {
178             log_error( "ERROR: __IMAGE_SUPPORT__ defined to value %d even though images aren't supported", (int)results[ 0 ] );
179             return -1;
180         }
181     }
182 
183     // __ENDIAN_LITTLE__ is similar to __IMAGE_SUPPORT__: 1 if it's true, undefined if it isn't
184     cl_bool deviceIsLittleEndian;
185     error = clGetDeviceInfo( deviceID, CL_DEVICE_ENDIAN_LITTLE, sizeof( deviceIsLittleEndian ), &deviceIsLittleEndian, NULL );
186     test_error( error, "Unable to get endian property of device to validate against" );
187 
188     if( deviceIsLittleEndian )
189     {
190         if( results[ 1 ] == 0xf00baa )
191         {
192             log_error( "ERROR: __ENDIAN_LITTLE__ undefined even though the device is little endian\n" );
193             return -1;
194         }
195         else if( results[ 1 ] != 1 )
196         {
197             log_error( "ERROR: __ENDIAN_LITTLE__ defined, but to the wrong value (defined as %d, spec states it should be 1)\n", (int)results[ 1 ] );
198             return -1;
199         }
200     }
201     else
202     {
203         if( results[ 1 ] != 0xf00baa )
204         {
205             log_error( "ERROR: __ENDIAN_LITTLE__ defined to value %d even though the device is not little endian (should be undefined per spec)", (int)results[ 1 ] );
206             return -1;
207         }
208     }
209 
210     // __OPENCL_VERSION__
211     if( results[ 2 ] == 0xf00baa )
212     {
213         log_error( "ERROR: Kernel preprocessor __OPENCL_VERSION__ undefined!" );
214         return -1;
215     }
216 
217     // The OpenCL version reported by the macro reports the feature level supported by the compiler. Since
218     // this doesn't directly match any property we can query, we just check to see if it's a sane value
219     auto device_cl_version = get_device_cl_version(deviceID);
220     int device_cl_version_int = device_cl_version.to_int() * 10;
221     if ((results[2] < 100) || (results[2] > device_cl_version_int))
222     {
223         log_error("ERROR: Kernel preprocessor __OPENCL_VERSION__ does not make "
224                   "sense w.r.t. device's version string! "
225                   "(preprocessor states %d, CL_DEVICE_VERSION is %d (%s))\n",
226                   results[2], device_cl_version_int,
227                   device_cl_version.to_string().c_str());
228         return -1;
229     }
230 
231     // __OPENCL_C_VERSION__
232     if( results[ 3 ] == 0xf00baa )
233     {
234         log_error( "ERROR: Kernel preprocessor __OPENCL_C_VERSION__ undefined!\n" );
235         return -1;
236     }
237 
238     // The OpenCL C version reported by the macro reports the OpenCL C version
239     // specified to the compiler. We need to see whether it is supported.
240     int cl_c_major_version = results[3] / 100;
241     int cl_c_minor_version = (results[3] / 10) % 10;
242     if ((results[3] < 100)
243         || (!device_supports_cl_c_version(
244             deviceID, Version{ cl_c_major_version, cl_c_minor_version })))
245     {
246         auto device_version = get_device_cl_c_version(deviceID);
247         log_error(
248             "ERROR: Kernel preprocessor __OPENCL_C_VERSION__ does not make "
249             "sense w.r.t. device's version string! "
250             "(preprocessor states %d, CL_DEVICE_OPENCL_C_VERSION is %d (%s))\n",
251             results[3], device_version.to_int() * 10,
252             device_version.to_string().c_str());
253         log_error("This means that CL_DEVICE_OPENCL_C_VERSION < "
254                   "__OPENCL_C_VERSION__");
255         if (device_cl_version >= Version{ 3, 0 })
256         {
257             log_error(", and __OPENCL_C_VERSION__ does not appear in "
258                       "CL_DEVICE_OPENCL_C_ALL_VERSIONS");
259         }
260         log_error("\n");
261         return -1;
262     }
263 
264     // __LINE__
265     if( results[ 4 ] == 0xf00baa )
266     {
267         log_error( "ERROR: Kernel preprocessor __LINE__ undefined!" );
268         return -1;
269     }
270 
271     // This is fun--we get to search for where __LINE__ actually is so we know what line it should define to!
272     // Note: it shows up twice, once for the #ifdef, and the other for the actual result output
273     const char *linePtr = strstr( preprocessor_test, "__LINE__" );
274     if( linePtr == NULL )
275     {
276         log_error( "ERROR: Nonsensical NULL pointer encountered!" );
277         return -2;
278     }
279     linePtr = strstr( linePtr + strlen( "__LINE__" ), "__LINE__" );
280     if( linePtr == NULL )
281     {
282         log_error( "ERROR: Nonsensical NULL pointer encountered!" );
283         return -2;
284     }
285 
286     // Now count how many carriage returns are before the string
287     const char *retPtr = strchr( preprocessor_test, '\n' );
288     int retCount = 1;
289     for( ; ( retPtr < linePtr ) && ( retPtr != NULL ); retPtr = strchr( retPtr + 1, '\n' ) )
290         retCount++;
291 
292     if( retCount != results[ 4 ] )
293     {
294         log_error( "ERROR: Kernel preprocessor __LINE__ does not expand to the actual line number! (expanded to %d, but was on line %d)\n",
295                   results[ 4 ], retCount );
296         return -1;
297     }
298 
299 #if 0 // Removed by Affie's request 2/24
300     // __FAST_RELAXED_MATH__
301     // Since create_single_kernel_helper does NOT define -cl-fast-relaxed-math, this should be undefined
302     if( results[ 5 ] != 0xf00baa )
303     {
304         log_error( "ERROR: Kernel preprocessor __FAST_RELAXED_MATH__ defined even though build option was not used (should be undefined)\n" );
305         return -1;
306     }
307 #endif
308 
309     // __kernel_exec
310     // We can ONLY check to verify that it is defined
311     if( results[ 6 ] == 0xf00baa )
312     {
313         log_error( "ERROR: Kernel preprocessor __kernel_exec must be defined\n" );
314         return -1;
315     }
316 
317     //// String preprocessors
318 
319     // Since we provided the program directly, __FILE__ should compile to "<program source>".
320     if( fileString[ 0 ] == 0 )
321     {
322         log_error( "ERROR: Kernel preprocessor __FILE__ undefined!\n" );
323         return -1;
324     }
325     else if( strncmp( (char *)fileString, __FILE__, 512 ) != 0 )
326     {
327         log_info( "WARNING: __FILE__ defined, but to an unexpected value (%s)\n\tShould be: \"%s\"", fileString, __FILE__ );
328         return -1;
329     }
330 
331 
332 #if 0 // Removed by Affie's request 2/24
333     // One more try through: try with -cl-fast-relaxed-math to make sure the appropriate preprocessor gets defined
334     clProgramWrapper programB = clCreateProgramWithSource( context, 1, preprocessor_test, NULL, &error );
335     test_error( error, "Unable to create test program" );
336 
337     // Try compiling
338     error = clBuildProgram( programB, 1, &deviceID, "-cl-fast-relaxed-math", NULL, NULL );
339     test_error( error, "Unable to build program" );
340 
341     // Create a kernel again to run against
342     clKernelWrapper kernelB = clCreateKernel( programB, "test", &error );
343     test_error( error, "Unable to create testing kernel" );
344 
345     // Set up and run
346     for( int i = 0; i < 3; i++ )
347     {
348         error = clSetKernelArg( kernelB, i, sizeof( streams[i] ), &streams[i] );
349         test_error( error, "Unable to set indexed kernel arguments" );
350     }
351 
352     error = clEnqueueNDRangeKernel( queue, kernelB, 1, NULL, threads, NULL, 0, NULL, NULL );
353     test_error( error, "Kernel execution failed" );
354 
355     // Only need the one read
356     error = clEnqueueReadBuffer( queue, streams[0], CL_TRUE, 0, sizeof(results), results, 0, NULL, NULL );
357     test_error( error, "Unable to get result data" );
358 
359     // We only need to check the one result this time
360     if( results[ 5 ] == 0xf00baa )
361     {
362         log_error( "ERROR: Kernel preprocessor __FAST_RELAXED_MATH__ not defined!\n" );
363         return -1;
364     }
365     else if( results[ 5 ] != 1 )
366     {
367         log_error( "ERROR: Kernel preprocessor __FAST_RELAXED_MATH__ not defined to 1 (was %d)\n", results[ 5 ] );
368         return -1;
369     }
370 #endif
371 
372     return 0;
373 }
374 
375