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 "harness/os_helpers.h"
18
19 const char *preprocessor_test_kernel[] = {
20 "__kernel void sample_test(__global int *dst)\n"
21 "{\n"
22 " dst[0] = TEST_MACRO;\n"
23 "\n"
24 "}\n" };
25
26 const char *preprocessor_existence_test_kernel[] = {
27 "__kernel void sample_test(__global int *dst)\n"
28 "{\n"
29 "#ifdef TEST_MACRO\n"
30 " dst[0] = 42;\n"
31 "#else\n"
32 " dst[0] = 24;\n"
33 "#endif\n"
34 "\n"
35 "}\n" };
36
37 const char *include_test_kernel[] = {
38 "#include \"./testIncludeFile.h\"\n"
39 "__kernel void sample_test(__global int *dst)\n"
40 "{\n"
41 " dst[0] = HEADER_FOUND;\n"
42 "\n"
43 "}\n" };
44
45 const char *options_test_kernel[] = {
46 "__kernel void sample_test(__global float *src, __global int *dst)\n"
47 "{\n"
48 " size_t tid = get_global_id(0);\n"
49 " dst[tid] = (int)src[tid];\n"
50 "}\n"
51 };
52
53 const char *optimization_options[] = {
54 "-cl-single-precision-constant",
55 "-cl-denorms-are-zero",
56 "-cl-opt-disable",
57 "-cl-mad-enable",
58 "-cl-no-signed-zeros",
59 "-cl-unsafe-math-optimizations",
60 "-cl-finite-math-only",
61 "-cl-fast-relaxed-math",
62 "-w",
63 "-Werror",
64 };
65
get_result_from_program(cl_context context,cl_command_queue queue,cl_program program,cl_int * outValue)66 cl_int get_result_from_program( cl_context context, cl_command_queue queue, cl_program program, cl_int *outValue )
67 {
68 cl_int error;
69 clKernelWrapper kernel = clCreateKernel( program, "sample_test", &error );
70 test_error( error, "Unable to create kernel from program" );
71
72 clMemWrapper outStream;
73 outStream = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int), NULL,
74 &error);
75 test_error( error, "Unable to create test buffer" );
76
77 error = clSetKernelArg( kernel, 0, sizeof( outStream ), &outStream );
78 test_error( error, "Unable to set kernel argument" );
79
80 size_t threads[1] = { 1 };
81
82 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
83 test_error( error, "Unable to execute test kernel" );
84
85 error = clEnqueueReadBuffer( queue, outStream, true, 0, sizeof( cl_int ), outValue, 0, NULL, NULL );
86 test_error( error, "Unable to read output array!" );
87
88 return CL_SUCCESS;
89 }
90
test_options_build_optimizations(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)91 int test_options_build_optimizations(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
92 {
93 int error;
94 cl_build_status status;
95
96 for(size_t i = 0; i < sizeof(optimization_options) / (sizeof(char*)); i++) {
97
98 clProgramWrapper program;
99 error = create_single_kernel_helper_create_program(context, &program, 1, options_test_kernel, optimization_options[i]);
100 if( program == NULL || error != CL_SUCCESS )
101 {
102 log_error( "ERROR: Unable to create reference program!\n" );
103 return -1;
104 }
105
106 /* Build with the macro defined */
107 log_info("Testing optimization option '%s'\n", optimization_options[i]);
108 error = clBuildProgram( program, 1, &deviceID, optimization_options[i], NULL, NULL );
109 test_error( error, "Test program did not properly build" );
110
111 error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
112 test_error( error, "Unable to get program build status" );
113
114 if( (int)status != CL_BUILD_SUCCESS )
115 {
116 log_info("Building with optimization option '%s' failed to compile!\n", optimization_options[i]);
117 print_error( error, "Failed to build with optimization defined")
118 return -1;
119 }
120 }
121 return 0;
122 }
123
test_options_build_macro(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)124 int test_options_build_macro(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
125 {
126 int error;
127 clProgramWrapper program;
128 cl_build_status status;
129
130
131 program = clCreateProgramWithSource( context, 1, preprocessor_test_kernel, NULL, &error );
132 if( program == NULL || error != CL_SUCCESS )
133 {
134 log_error( "ERROR: Unable to create reference program!\n" );
135 return -1;
136 }
137
138 /* Build with the macro defined */
139 error = clBuildProgram( program, 1, &deviceID, "-DTEST_MACRO=1 ", NULL, NULL );
140 test_error( error, "Test program did not properly build" );
141
142 error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
143 test_error( error, "Unable to get program build status" );
144
145 if( (int)status != CL_BUILD_SUCCESS )
146 {
147 print_error( error, "Failed to build with macro defined" );
148 return -1;
149 }
150
151
152 // Go ahead and run the program to verify results
153 cl_int firstResult, secondResult;
154
155 error = get_result_from_program( context, queue, program, &firstResult );
156 test_error( error, "Unable to get result from first program" );
157
158 if( firstResult != 1 )
159 {
160 log_error( "ERROR: Result from first program did not validate! (Expected 1, got %d)\n", firstResult );
161 return -1;
162 }
163
164 // Rebuild with a different value for the define macro, to make sure caching behaves properly
165 error = clBuildProgram( program, 1, &deviceID, "-DTEST_MACRO=5 ", NULL, NULL );
166 test_error( error, "Test program did not properly rebuild" );
167
168 error = get_result_from_program( context, queue, program, &secondResult );
169 test_error( error, "Unable to get result from second program" );
170
171 if( secondResult != 5 )
172 {
173 if( secondResult == firstResult )
174 log_error( "ERROR: Program result did not change with device macro change (program was not recompiled)!\n" );
175 else
176 log_error( "ERROR: Result from second program did not validate! (Expected 5, got %d)\n", secondResult );
177 return -1;
178 }
179
180 return 0;
181 }
182
test_options_build_macro_existence(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)183 int test_options_build_macro_existence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
184 {
185 int error;
186 clProgramWrapper program;
187
188
189 // In this case, the program should still run without the macro, but it should give a different result
190 program = clCreateProgramWithSource( context, 1, preprocessor_existence_test_kernel, NULL, &error );
191 if( program == NULL || error != CL_SUCCESS )
192 {
193 log_error( "ERROR: Unable to create reference program!\n" );
194 return -1;
195 }
196
197 /* Build without the macro defined */
198 error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
199 test_error( error, "Test program did not properly build" );
200
201 // Go ahead and run the program to verify results
202 cl_int firstResult, secondResult;
203
204 error = get_result_from_program( context, queue, program, &firstResult );
205 test_error( error, "Unable to get result from first program" );
206
207 if( firstResult != 24 )
208 {
209 log_error( "ERROR: Result from first program did not validate! (Expected 24, got %d)\n", firstResult );
210 return -1;
211 }
212
213 // Now compile again with the macro defined and verify a change in results
214 error = clBuildProgram( program, 1, &deviceID, "-DTEST_MACRO", NULL, NULL );
215 test_error( error, "Test program did not properly build" );
216
217 error = get_result_from_program( context, queue, program, &secondResult );
218 test_error( error, "Unable to get result from second program" );
219
220 if( secondResult != 42 )
221 {
222 if( secondResult == firstResult )
223 log_error( "ERROR: Program result did not change with device macro addition (program was not recompiled)!\n" );
224 else
225 log_error( "ERROR: Result from second program did not validate! (Expected 42, got %d)\n", secondResult );
226 return -1;
227 }
228
229 return 0;
230 }
231
test_options_include_directory(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)232 int test_options_include_directory(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
233 {
234 int error;
235
236 std::string sep = dir_sep();
237 std::string path = exe_dir(); // Directory where test executable is located.
238 std::string include_dir;
239
240 clProgramWrapper program;
241 cl_build_status status;
242
243 /* Try compiling the program first without the directory included Should fail. */
244 program = clCreateProgramWithSource( context, 1, include_test_kernel, NULL, &error );
245 if( program == NULL || error != CL_SUCCESS )
246 {
247 log_error( "ERROR: Unable to create reference program!\n" );
248 return -1;
249 }
250
251 /* Build with the include directory defined */
252 include_dir = "-I " + path + sep + "includeTestDirectory";
253
254 // log_info("%s\n", include_dir);
255 error = clBuildProgram( program, 1, &deviceID, include_dir.c_str(), NULL, NULL );
256 test_error( error, "Test program did not properly build" );
257
258 error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
259 test_error( error, "Unable to get program build status" );
260
261 if( (int)status != CL_BUILD_SUCCESS )
262 {
263 print_error( error, "Failed to build with include directory" );
264 return -1;
265 }
266
267 // Go ahead and run the program to verify results
268 cl_int firstResult, secondResult;
269
270 error = get_result_from_program( context, queue, program, &firstResult );
271 test_error( error, "Unable to get result from first program" );
272
273 if( firstResult != 12 )
274 {
275 log_error( "ERROR: Result from first program did not validate! (Expected 12, got %d)\n", firstResult );
276 return -1;
277 }
278
279 // Rebuild with a different include directory
280 include_dir = "-I " + path + sep + "secondIncludeTestDirectory";
281 error = clBuildProgram( program, 1, &deviceID, include_dir.c_str(), NULL, NULL );
282 test_error( error, "Test program did not properly rebuild" );
283
284 error = get_result_from_program( context, queue, program, &secondResult );
285 test_error( error, "Unable to get result from second program" );
286
287 if( secondResult != 42 )
288 {
289 if( secondResult == firstResult )
290 log_error( "ERROR: Program result did not change with include path change (program was not recompiled)!\n" );
291 else
292 log_error( "ERROR: Result from second program did not validate! (Expected 42, got %d)\n", secondResult );
293 return -1;
294 }
295
296 return 0;
297 }
298
299 const char *denorm_test_kernel[] = {
300 "__kernel void sample_test( float src1, float src2, __global float *dst)\n"
301 "{\n"
302 " dst[ 0 ] = src1 + src2;\n"
303 "\n"
304 "}\n" };
305
get_float_result_from_program(cl_context context,cl_command_queue queue,cl_program program,cl_float inA,cl_float inB,cl_float * outValue)306 cl_int get_float_result_from_program( cl_context context, cl_command_queue queue, cl_program program, cl_float inA, cl_float inB, cl_float *outValue )
307 {
308 cl_int error;
309
310 clKernelWrapper kernel = clCreateKernel( program, "sample_test", &error );
311 test_error( error, "Unable to create kernel from program" );
312
313 clMemWrapper outStream = clCreateBuffer(context, CL_MEM_READ_WRITE,
314 sizeof(cl_float), NULL, &error);
315 test_error( error, "Unable to create test buffer" );
316
317 error = clSetKernelArg( kernel, 0, sizeof( cl_float ), &inA );
318 test_error( error, "Unable to set kernel argument" );
319
320 error = clSetKernelArg( kernel, 1, sizeof( cl_float ), &inB );
321 test_error( error, "Unable to set kernel argument" );
322
323 error = clSetKernelArg( kernel, 2, sizeof( outStream ), &outStream );
324 test_error( error, "Unable to set kernel argument" );
325
326 size_t threads[1] = { 1 };
327
328 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
329 test_error( error, "Unable to execute test kernel" );
330
331 error = clEnqueueReadBuffer( queue, outStream, true, 0, sizeof( cl_float ), outValue, 0, NULL, NULL );
332 test_error( error, "Unable to read output array!" );
333
334 return CL_SUCCESS;
335 }
336
test_options_denorm_cache(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)337 int test_options_denorm_cache(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
338 {
339 int error;
340
341 clProgramWrapper program;
342 cl_build_status status;
343
344
345 // If denorms aren't even supported, testing this flag is pointless
346 cl_device_fp_config floatCaps = 0;
347 error = clGetDeviceInfo( deviceID, CL_DEVICE_SINGLE_FP_CONFIG, sizeof(floatCaps), &floatCaps, NULL);
348 test_error( error, "Unable to get device FP config" );
349 if( ( floatCaps & CL_FP_DENORM ) == 0 )
350 {
351 log_info( "Device does not support denormalized single-precision floats; skipping test.\n" );
352 return 0;
353 }
354
355 program = clCreateProgramWithSource( context, 1, denorm_test_kernel, NULL, &error );
356 test_error( error, "Unable to create test program" );
357
358 // Build first WITH the denorm flush flag
359 error = clBuildProgram( program, 1, &deviceID, "-cl-denorms-are-zero", NULL, NULL );
360 test_error( error, "Test program did not properly build" );
361
362 error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL );
363 test_error( error, "Unable to get program build status" );
364
365 if( (int)status != CL_BUILD_SUCCESS )
366 {
367 print_error( error, "Failed to build with include directory" );
368 return -1;
369 }
370
371 // Note: the following in floating point is a subnormal number, thus adding two of them together
372 // should give us a subnormalized result. If denormals are flushed to zero, however, it'll give us zero instead
373 uint32_t intSubnormal = 0x00000001;
374 cl_float *input = (cl_float *)&intSubnormal;
375 cl_float firstResult, secondResult;
376
377 error = get_float_result_from_program( context, queue, program, *input, *input, &firstResult );
378 test_error( error, "Unable to get result from first program" );
379
380 // Note: since -cl-denorms-are-zero is a HINT, not a requirement, the result we got could
381 // either be subnormal (hint ignored) or zero (hint respected). Since either is technically
382 // valid, there isn't anything we can to do validate results for now
383
384 // Rebuild without flushing flag set
385 error = clBuildProgram( program, 1, &deviceID, NULL, NULL, NULL );
386 test_error( error, "Test program did not properly rebuild" );
387
388 error = get_float_result_from_program( context, queue, program, *input, *input, &secondResult );
389 test_error( error, "Unable to get result from second program" );
390
391 // Now, there are three possiblities here:
392 // 1. The denorms-are-zero hint is not respected, in which case the first and second result will be identical
393 // 2. The hint is respected, and the program was properly rebuilt, in which case the first result will be zero and the second non-zero
394 // 3. The hint is respected, but the program was not properly rebuilt, in which case both results will be zero
395 // 3 is the only error condition we need to look for
396 uint32_t *fPtr = (uint32_t *)&firstResult;
397 uint32_t *sPtr = (uint32_t *)&secondResult;
398
399 if( ( *fPtr == 0 ) && ( *sPtr == 0 ) )
400 {
401 log_error( "ERROR: Program result didn't change when -cl-denorms-are-zero flag was removed.\n"
402 "First result (should be zero): 0x%08x, Second result (should be non-zero): 0x%08x\n",
403 *fPtr, *sPtr );
404 return -1;
405 }
406
407 return 0;
408 }
409
410