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