test_conformance/basic/test_preprocessors.cpp

//
// 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 <ctype.h>

// 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 "<program source>".
    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;
}