// // Copyright (c) 2017 The Khronos Group Inc. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at // // http://www.apache.org/licenses/LICENSE-2.0 // // Unless required by applicable law or agreed to in writing, software // distributed under the License is distributed on an "AS IS" BASIS, // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // See the License for the specific language governing permissions and // limitations under the License. // #include "procs.h" #include // Test __FILE__, __LINE__, __OPENCL_VERSION__, __OPENCL_C_VERSION__, __ENDIAN_LITTLE__, __ROUNDING_MODE__, __IMAGE_SUPPORT__, __FAST_RELAXED_MATH__ // __kernel_exec const char *preprocessor_test = { "#line 2 \"%s\"\n" "__kernel void test( __global int *results, __global char *outFileString, __global char *outRoundingString )\n" "{\n" // Integer preprocessor macros "#ifdef __IMAGE_SUPPORT__\n" " results[0] = __IMAGE_SUPPORT__;\n" "#else\n" " results[0] = 0xf00baa;\n" "#endif\n" "#ifdef __ENDIAN_LITTLE__\n" " results[1] = __ENDIAN_LITTLE__;\n" "#else\n" " results[1] = 0xf00baa;\n" "#endif\n" "#ifdef __OPENCL_VERSION__\n" " results[2] = __OPENCL_VERSION__;\n" "#else\n" " results[2] = 0xf00baa;\n" "#endif\n" "#ifdef __OPENCL_C_VERSION__\n" " results[3] = __OPENCL_C_VERSION__;\n" "#else\n" " results[3] = 0xf00baa;\n" "#endif\n" "#ifdef __LINE__\n" " results[4] = __LINE__;\n" "#else\n" " results[4] = 0xf00baa;\n" "#endif\n" #if 0 // Removed by Affie's request 2/24 "#ifdef __FAST_RELAXED_MATH__\n" " results[5] = __FAST_RELAXED_MATH__;\n" "#else\n" " results[5] = 0xf00baa;\n" "#endif\n" #endif "#ifdef __kernel_exec\n" " results[6] = 1;\n" // By spec, we can only really evaluate that it is defined, not what it expands to "#else\n" " results[6] = 0xf00baa;\n" "#endif\n" // String preprocessor macros. Technically, there are strings in OpenCL, but not really. "#ifdef __FILE__\n" " int i;\n" " constant char *f = \"\" __FILE__;\n" " for( i = 0; f[ i ] != 0 && i < 512; i++ )\n" " outFileString[ i ] = f[ i ];\n" " outFileString[ i ] = 0;\n" "#else\n" " outFileString[ 0 ] = 0;\n" "#endif\n" "}\n" }; int test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { clProgramWrapper program; clKernelWrapper kernel; clMemWrapper streams[ 3 ]; int error; size_t threads[] = {1,1,1}; cl_int results[ 7 ]; cl_char fileString[ 512 ] = "", roundingString[ 128 ] = ""; char programSource[4096]; char curFileName[512]; char *programPtr = programSource; int i = 0; snprintf(curFileName, 512, "%s", __FILE__); #ifdef _WIN32 // Replace "\" with "\\" while(curFileName[i] != '\0') { if (curFileName[i] == '\\') { int j = i + 1; char prev = '\\'; while (curFileName[j - 1] != '\0') { char tmp = curFileName[j]; curFileName[j] = prev; prev = tmp; j++; } i++; } i++; } #endif sprintf(programSource,preprocessor_test,curFileName); // Create the kernel if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "test" ) != 0 ) { return -1; } /* Create some I/O streams */ streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(results), NULL, &error); test_error( error, "Creating test array failed" ); streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(fileString), NULL, &error); test_error( error, "Creating test array failed" ); streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(roundingString), NULL, &error); test_error( error, "Creating test array failed" ); // Set up and run for( int i = 0; i < 3; i++ ) { error = clSetKernelArg( kernel, i, sizeof( streams[i] ), &streams[i] ); test_error( error, "Unable to set indexed kernel arguments" ); } error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL ); test_error( error, "Kernel execution failed" ); error = clEnqueueReadBuffer( queue, streams[0], CL_TRUE, 0, sizeof(results), results, 0, NULL, NULL ); test_error( error, "Unable to get result data" ); error = clEnqueueReadBuffer( queue, streams[1], CL_TRUE, 0, sizeof(fileString), fileString, 0, NULL, NULL ); test_error( error, "Unable to get result data" ); error = clEnqueueReadBuffer( queue, streams[2], CL_TRUE, 0, sizeof(roundingString), roundingString, 0, NULL, NULL ); test_error( error, "Unable to get result data" ); /////// Check the integer results // We need to check these values against what we know is supported on the device if( checkForImageSupport( deviceID ) == 0 ) { // If images are supported, the constant should have been defined to the value 1 if( results[ 0 ] == 0xf00baa ) { log_error( "ERROR: __IMAGE_SUPPORT__ undefined even though images are supported\n" ); return -1; } else if( results[ 0 ] != 1 ) { log_error( "ERROR: __IMAGE_SUPPORT__ defined, but to the wrong value (defined as %d, spec states it should be 1)\n", (int)results[ 0 ] ); return -1; } } else { // If images aren't supported, the constant should be undefined if( results[ 0 ] != 0xf00baa ) { log_error( "ERROR: __IMAGE_SUPPORT__ defined to value %d even though images aren't supported", (int)results[ 0 ] ); return -1; } } // __ENDIAN_LITTLE__ is similar to __IMAGE_SUPPORT__: 1 if it's true, undefined if it isn't cl_bool deviceIsLittleEndian; error = clGetDeviceInfo( deviceID, CL_DEVICE_ENDIAN_LITTLE, sizeof( deviceIsLittleEndian ), &deviceIsLittleEndian, NULL ); test_error( error, "Unable to get endian property of device to validate against" ); if( deviceIsLittleEndian ) { if( results[ 1 ] == 0xf00baa ) { log_error( "ERROR: __ENDIAN_LITTLE__ undefined even though the device is little endian\n" ); return -1; } else if( results[ 1 ] != 1 ) { log_error( "ERROR: __ENDIAN_LITTLE__ defined, but to the wrong value (defined as %d, spec states it should be 1)\n", (int)results[ 1 ] ); return -1; } } else { if( results[ 1 ] != 0xf00baa ) { 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 ] ); return -1; } } // __OPENCL_VERSION__ if( results[ 2 ] == 0xf00baa ) { log_error( "ERROR: Kernel preprocessor __OPENCL_VERSION__ undefined!" ); return -1; } // The OpenCL version reported by the macro reports the feature level supported by the compiler. Since // this doesn't directly match any property we can query, we just check to see if it's a sane value auto device_cl_version = get_device_cl_version(deviceID); int device_cl_version_int = device_cl_version.to_int() * 10; if ((results[2] < 100) || (results[2] > device_cl_version_int)) { log_error("ERROR: Kernel preprocessor __OPENCL_VERSION__ does not make " "sense w.r.t. device's version string! " "(preprocessor states %d, CL_DEVICE_VERSION is %d (%s))\n", results[2], device_cl_version_int, device_cl_version.to_string().c_str()); return -1; } // __OPENCL_C_VERSION__ if( results[ 3 ] == 0xf00baa ) { log_error( "ERROR: Kernel preprocessor __OPENCL_C_VERSION__ undefined!\n" ); return -1; } // The OpenCL C version reported by the macro reports the OpenCL C version // specified to the compiler. We need to see whether it is supported. int cl_c_major_version = results[3] / 100; int cl_c_minor_version = (results[3] / 10) % 10; if ((results[3] < 100) || (!device_supports_cl_c_version( deviceID, Version{ cl_c_major_version, cl_c_minor_version }))) { auto device_version = get_device_cl_c_version(deviceID); log_error( "ERROR: Kernel preprocessor __OPENCL_C_VERSION__ does not make " "sense w.r.t. device's version string! " "(preprocessor states %d, CL_DEVICE_OPENCL_C_VERSION is %d (%s))\n", results[3], device_version.to_int() * 10, device_version.to_string().c_str()); log_error("This means that CL_DEVICE_OPENCL_C_VERSION < " "__OPENCL_C_VERSION__"); if (device_cl_version >= Version{ 3, 0 }) { log_error(", and __OPENCL_C_VERSION__ does not appear in " "CL_DEVICE_OPENCL_C_ALL_VERSIONS"); } log_error("\n"); return -1; } // __LINE__ if( results[ 4 ] == 0xf00baa ) { log_error( "ERROR: Kernel preprocessor __LINE__ undefined!" ); return -1; } // This is fun--we get to search for where __LINE__ actually is so we know what line it should define to! // Note: it shows up twice, once for the #ifdef, and the other for the actual result output const char *linePtr = strstr( preprocessor_test, "__LINE__" ); if( linePtr == NULL ) { log_error( "ERROR: Nonsensical NULL pointer encountered!" ); return -2; } linePtr = strstr( linePtr + strlen( "__LINE__" ), "__LINE__" ); if( linePtr == NULL ) { log_error( "ERROR: Nonsensical NULL pointer encountered!" ); return -2; } // Now count how many carriage returns are before the string const char *retPtr = strchr( preprocessor_test, '\n' ); int retCount = 1; for( ; ( retPtr < linePtr ) && ( retPtr != NULL ); retPtr = strchr( retPtr + 1, '\n' ) ) retCount++; if( retCount != results[ 4 ] ) { log_error( "ERROR: Kernel preprocessor __LINE__ does not expand to the actual line number! (expanded to %d, but was on line %d)\n", results[ 4 ], retCount ); return -1; } #if 0 // Removed by Affie's request 2/24 // __FAST_RELAXED_MATH__ // Since create_single_kernel_helper does NOT define -cl-fast-relaxed-math, this should be undefined if( results[ 5 ] != 0xf00baa ) { log_error( "ERROR: Kernel preprocessor __FAST_RELAXED_MATH__ defined even though build option was not used (should be undefined)\n" ); return -1; } #endif // __kernel_exec // We can ONLY check to verify that it is defined if( results[ 6 ] == 0xf00baa ) { log_error( "ERROR: Kernel preprocessor __kernel_exec must be defined\n" ); return -1; } //// String preprocessors // Since we provided the program directly, __FILE__ should compile to "". if( fileString[ 0 ] == 0 ) { log_error( "ERROR: Kernel preprocessor __FILE__ undefined!\n" ); return -1; } else if( strncmp( (char *)fileString, __FILE__, 512 ) != 0 ) { log_info( "WARNING: __FILE__ defined, but to an unexpected value (%s)\n\tShould be: \"%s\"", fileString, __FILE__ ); return -1; } #if 0 // Removed by Affie's request 2/24 // One more try through: try with -cl-fast-relaxed-math to make sure the appropriate preprocessor gets defined clProgramWrapper programB = clCreateProgramWithSource( context, 1, preprocessor_test, NULL, &error ); test_error( error, "Unable to create test program" ); // Try compiling error = clBuildProgram( programB, 1, &deviceID, "-cl-fast-relaxed-math", NULL, NULL ); test_error( error, "Unable to build program" ); // Create a kernel again to run against clKernelWrapper kernelB = clCreateKernel( programB, "test", &error ); test_error( error, "Unable to create testing kernel" ); // Set up and run for( int i = 0; i < 3; i++ ) { error = clSetKernelArg( kernelB, i, sizeof( streams[i] ), &streams[i] ); test_error( error, "Unable to set indexed kernel arguments" ); } error = clEnqueueNDRangeKernel( queue, kernelB, 1, NULL, threads, NULL, 0, NULL, NULL ); test_error( error, "Kernel execution failed" ); // Only need the one read error = clEnqueueReadBuffer( queue, streams[0], CL_TRUE, 0, sizeof(results), results, 0, NULL, NULL ); test_error( error, "Unable to get result data" ); // We only need to check the one result this time if( results[ 5 ] == 0xf00baa ) { log_error( "ERROR: Kernel preprocessor __FAST_RELAXED_MATH__ not defined!\n" ); return -1; } else if( results[ 5 ] != 1 ) { log_error( "ERROR: Kernel preprocessor __FAST_RELAXED_MATH__ not defined to 1 (was %d)\n", results[ 5 ] ); return -1; } #endif return 0; }