//
// 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 "harness/compat.h"
#include "harness/rounding_mode.h"
#include "harness/ThreadPool.h"
#include "harness/testHarness.h"
#include "harness/kernelHelpers.h"
#include "harness/parseParameters.h"
#if defined(__APPLE__)
#include <sys/sysctl.h>
#endif

#if defined( __linux__ )
#include <unistd.h>
#include <sys/syscall.h>
#include <linux/sysctl.h>
#endif
#if defined(__linux__)
#include <sys/param.h>
#include <libgen.h>
#endif

#include "mingw_compat.h"
#if defined(__MINGW32__)
#include <sys/param.h>
#endif

#include <stdarg.h>
#include <stdio.h>
#include <string.h>
#if !defined(_WIN32)
#include <libgen.h>
#include <sys/mman.h>
#endif
#include <time.h>

#include "Sleep.h"
#include "basic_test_conversions.h"

#if (defined(_WIN32) && defined (_MSC_VER))
// need for _controlfp_s and rouinding modes in RoundingMode
#include "harness/testHarness.h"
#endif

#pragma mark -
#pragma mark globals

#define BUFFER_SIZE     (1024*1024)
#define kPageSize       4096
#define EMBEDDED_REDUCTION_FACTOR 16
#define PERF_LOOP_COUNT 100

#define      kCallStyleCount (kVectorSizeCount + 1 /* for implicit scalar */)

#if (defined(__arm__) || defined(__aarch64__)) && defined(__GNUC__)
#include "fplib.h"
    extern bool            qcom_sat;
    extern roundingMode    qcom_rm;
#endif

const char **   argList = NULL;
int             argCount = 0;
cl_context      gContext = NULL;
cl_command_queue      gQueue = NULL;
char            appName[64] = "ctest";
int             gStartTestNumber = -1;
int             gEndTestNumber = 0;
#if defined( __APPLE__ )
int             gTimeResults = 1;
#else
int             gTimeResults = 0;
#endif
int             gReportAverageTimes = 0;
void            *gIn = NULL;
void            *gRef = NULL;
void        *gAllowZ = NULL;
void            *gOut[ kCallStyleCount ] = { NULL };
cl_mem          gInBuffer;
cl_mem          gOutBuffers[ kCallStyleCount ];
size_t          gComputeDevices = 0;
uint32_t        gDeviceFrequency = 0;
int             gWimpyMode = 0;
int             gWimpyReductionFactor = 128;
int             gSkipTesting = 0;
int             gForceFTZ = 0;
int             gMultithread = 1;
int             gIsRTZ = 0;
uint32_t        gSimdSize = 1;
int             gHasDouble = 0;
int             gTestDouble = 1;
const char *    sizeNames[] = { "", "", "2", "3", "4", "8", "16" };
const int       vectorSizes[] = { 1, 1, 2, 3, 4, 8, 16 };
int             gMinVectorSize = 0;
int             gMaxVectorSize = sizeof(vectorSizes) / sizeof( vectorSizes[0] );
static MTdata   gMTdata;

#pragma mark -
#pragma mark Declarations

static int ParseArgs( int argc, const char **argv );
static void PrintUsage( void );
test_status InitCL( cl_device_id device );
static int GetTestCase( const char *name, Type *outType, Type *inType, SaturationMode *sat, RoundingMode *round );
static int DoTest( cl_device_id device, Type outType, Type inType, SaturationMode sat, RoundingMode round, MTdata d );
static cl_program   MakeProgram( Type outType, Type inType, SaturationMode sat, RoundingMode round, int vectorSize, cl_kernel *outKernel );
static int RunKernel( cl_kernel kernel, void *inBuf, void *outBuf, size_t blockCount );

void *FlushToZero( void );
void UnFlushToZero( void *);

static cl_program CreateImplicitConvertProgram( Type outType, Type inType, SaturationMode sat, RoundingMode round, int vectorSize, char testName[256], cl_int *error );
static cl_program CreateStandardProgram( Type outType, Type inType, SaturationMode sat, RoundingMode round, int vectorSize, char testName[256], cl_int *error );


// Windows (since long double got deprecated) sets the x87 to 53-bit precision
// (that's x87 default state).  This causes problems with the tests that
// convert long and ulong to float and double or otherwise deal with values
// that need more precision than 53-bit. So, set the x87 to 64-bit precision.
static inline void Force64BitFPUPrecision(void)
{
#if __MINGW32__
    // The usual method is to use _controlfp as follows:
    //     #include <float.h>
    //     _controlfp(_PC_64, _MCW_PC);
    //
    // _controlfp is available on MinGW32 but not on MinGW64. Instead of having
    // divergent code just use inline assembly which works for both.
    unsigned short int orig_cw = 0;
    unsigned short int new_cw = 0;
    __asm__ __volatile__ ("fstcw %0":"=m" (orig_cw));
    new_cw = orig_cw | 0x0300;   // set precision to 64-bit
    __asm__ __volatile__ ("fldcw  %0"::"m" (new_cw));
#else
    /* Implement for other platforms if needed */
#endif
}

int test_conversions( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
{
    int error, i, testNumber = -1;
    int startMinVectorSize = gMinVectorSize;
    Type inType, outType;
    RoundingMode round;
    SaturationMode sat;

    if( argCount )
    {
        for( i = 0; i < argCount; i++ )
        {
            if( GetTestCase( argList[i], &outType, &inType, &sat, &round ) )
            {
                vlog_error( "\n\t\t**** ERROR:  Unable to parse function name %s.  Skipping....  *****\n\n", argList[i] );
                continue;
            }

            // skip double if we don't have it
            if( !gTestDouble && (inType == kdouble || outType == kdouble ) )
            {
                if( gHasDouble )
                {
                    vlog_error( "\t *** convert_%sn%s%s( %sn ) FAILED ** \n", gTypeNames[ outType ], gSaturationNames[ sat ], gRoundingModeNames[round], gTypeNames[inType] );
                    vlog( "\t\tcl_khr_fp64 enabled, but double testing turned off.\n" );
                }

                continue;
            }

            // skip longs on embedded
            if( !gHasLong && (inType == klong || outType == klong || inType == kulong || outType == kulong) )
            {
                continue;
            }

            // Skip the implicit converts if the rounding mode is not default or test is saturated
            if( 0 == startMinVectorSize )
            {
                if( sat || round != kDefaultRoundingMode )
                    gMinVectorSize = 1;
                else
                    gMinVectorSize = 0;
            }

            if( ( error = DoTest( device, outType, inType, sat, round, gMTdata ) ) )
            {
                vlog_error( "\t *** convert_%sn%s%s( %sn ) FAILED ** \n", gTypeNames[outType], gSaturationNames[sat], gRoundingModeNames[round], gTypeNames[inType] );
            }
        }
    }
    else
    {
        for( outType = (Type)0; outType < kTypeCount; outType = (Type)(outType+1) )
        {
            for( inType = (Type)0; inType < kTypeCount; inType = (Type)(inType+1) )
            {
                // skip longs on embedded
                if( !gHasLong && (inType == klong || outType == klong || inType == kulong || outType == kulong) )
                {
                    continue;
                }

                for( sat = (SaturationMode)0; sat < kSaturationModeCount; sat = (SaturationMode)(sat+1) )
                {
                    //skip illegal saturated conversions to float type
                    if( kSaturated == sat && ( outType == kfloat || outType == kdouble ) )
                    {
                        continue;
                    }

                    for( round = (RoundingMode)0; round < kRoundingModeCount; round = (RoundingMode)(round+1) )
                    {
                        if( ++testNumber < gStartTestNumber )
                        {
                            //     vlog( "%d) skipping convert_%sn%s%s( %sn )\n", testNumber, gTypeNames[ outType ], gSaturationNames[ sat ], gRoundingModeNames[round], gTypeNames[inType] );
                            continue;
                        }
                        else
                        {
                            if( gEndTestNumber > 0 && testNumber >= gEndTestNumber  )
                            {
                                goto exit;
                            }
                        }

                        vlog( "%d) Testing convert_%sn%s%s( %sn ):\n", testNumber, gTypeNames[ outType ], gSaturationNames[ sat ], gRoundingModeNames[round], gTypeNames[inType] );

                        // skip double if we don't have it
                        if( ! gTestDouble && (inType == kdouble || outType == kdouble ) )
                        {
                            if( gHasDouble )
                            {
                                vlog_error( "\t *** %d) convert_%sn%s%s( %sn ) FAILED ** \n", testNumber, gTypeNames[ outType ], gSaturationNames[ sat ], gRoundingModeNames[round], gTypeNames[inType] );
                                vlog( "\t\tcl_khr_fp64 enabled, but double testing turned off.\n" );
                            }
                            continue;
                        }

                        // Skip the implicit converts if the rounding mode is not default or test is saturated
                        if( 0 == startMinVectorSize )
                        {
                            if( sat || round != kDefaultRoundingMode )
                                gMinVectorSize = 1;
                            else
                                gMinVectorSize = 0;
                        }

                        if( ( error = DoTest( device, outType, inType, sat, round, gMTdata ) ) )
                        {
                            vlog_error( "\t *** %d) convert_%sn%s%s( %sn ) FAILED ** \n", testNumber, gTypeNames[outType], gSaturationNames[sat], gRoundingModeNames[round], gTypeNames[inType] );
                        }
                    }
                }
            }
        }
    }

exit:
    return gFailCount;
}

test_definition test_list[] = {
    ADD_TEST( conversions ),
};

const int test_num = ARRAY_SIZE( test_list );

#pragma mark -

int main (int argc, const char **argv )
{
    int error;
    cl_uint seed = (cl_uint) time( NULL );

    argc = parseCustomParam(argc, argv);
    if (argc == -1)
    {
        return 1;
    }

    if( (error = ParseArgs( argc, argv )) )
        return error;

    //Turn off sleep so our tests run to completion
    PreventSleep();
    atexit( ResumeSleep );

    if(!gMultithread)
        SetThreadCount(1);

#if defined(_MSC_VER) && defined(_M_IX86)
    // VS2005 (and probably others, since long double got deprecated) sets
    // the x87 to 53-bit precision. This causes problems with the tests
    // that convert long and ulong to float and double, since they deal
    // with values that need more precision than that. So, set the x87
    // to 64-bit precision.
    unsigned int ignored;
    _controlfp_s(&ignored, _PC_64, _MCW_PC);
#endif

    vlog( "===========================================================\n" );
    vlog( "Random seed: %u\n", seed );
    gMTdata = init_genrand( seed );

    const char* arg[] = {argv[0]};
    int ret = runTestHarnessWithCheck( 1, arg, test_num, test_list, true, 0, InitCL );

    free_mtdata( gMTdata );
    if (gQueue)
    {
        error = clFinish(gQueue);
        if (error) vlog_error("clFinish failed: %d\n", error);
    }

    clReleaseMemObject(gInBuffer);

    for( int i = 0; i < kCallStyleCount; i++ ) {
        clReleaseMemObject(gOutBuffers[i]);
    }
    clReleaseCommandQueue(gQueue);
    clReleaseContext(gContext);

    return ret;
}

#pragma mark -
#pragma mark setup

static int ParseArgs( int argc, const char **argv )
{
    int i;
    argList = (const char **)calloc( argc - 1, sizeof( char*) );
    argCount = 0;

    if( NULL == argList && argc > 1 )
        return -1;

#if (defined( __APPLE__ ) || defined(__linux__) || defined (__MINGW32__))
    { // Extract the app name
        char baseName[ MAXPATHLEN ];
        strncpy( baseName, argv[0], MAXPATHLEN );
        char *base = basename( baseName );
        if( NULL != base )
        {
            strncpy( appName, base, sizeof( appName )  );
            appName[ sizeof( appName ) -1 ] = '\0';
        }
    }
#elif defined (_WIN32)
    {
        char fname[_MAX_FNAME + _MAX_EXT + 1];
        char ext[_MAX_EXT];

        errno_t err = _splitpath_s( argv[0], NULL, 0, NULL, 0,
                                   fname, _MAX_FNAME, ext, _MAX_EXT );
        if (err == 0) { // no error
            strcat (fname, ext); //just cat them, size of frame can keep both
            strncpy (appName, fname, sizeof(appName));
            appName[ sizeof( appName ) -1 ] = '\0';
        }
    }
#endif

    vlog( "\n%s", appName );
    for( i = 1; i < argc; i++ )
    {
        const char *arg = argv[i];
        if( NULL == arg )
            break;

        vlog( "\t%s", arg );
        if( arg[0] == '-' )
        {
            arg++;
            while( *arg != '\0' )
            {
                switch( *arg )
                {
                    case 'd':
                        gTestDouble ^= 1;
                        break;
                    case 'l':
                        gSkipTesting ^= 1;
                        break;
                    case 'm':
                        gMultithread ^= 1;
                        break;
                    case 'w':
                        gWimpyMode ^= 1;
                        break;
                    case '[':
                        parseWimpyReductionFactor(arg, gWimpyReductionFactor);
                        break;
                    case 'z':
                        gForceFTZ ^= 1;
                        break;
                    case 't':
                        gTimeResults ^= 1;
                        break;
                    case 'a':
                        gReportAverageTimes ^= 1;
                        break;
                    case '1':
                        if( arg[1] == '6' )
                        {
                            gMinVectorSize = 6;
                            gMaxVectorSize = 7;
                            arg++;
                        }
                        else
                        {
                            gMinVectorSize = 0;
                            gMaxVectorSize = 2;
                        }
                        break;

                    case '2':
                        gMinVectorSize = 2;
                        gMaxVectorSize = 3;
                        break;

                    case '3':
                        gMinVectorSize = 3;
                        gMaxVectorSize = 4;
                        break;

                    case '4':
                        gMinVectorSize = 4;
                        gMaxVectorSize = 5;
                        break;

                    case '8':
                        gMinVectorSize = 5;
                        gMaxVectorSize = 6;
                        break;

                    default:
                        vlog( " <-- unknown flag: %c (0x%2.2x)\n)", *arg, *arg );
                        PrintUsage();
                        return -1;
                }
                arg++;
            }
        }
        else
        {
            char *t = NULL;
            long number = strtol( arg, &t, 0 );
            if( t != arg )
            {
                if( gStartTestNumber != -1 )
                    gEndTestNumber = gStartTestNumber + (int) number;
                else
                    gStartTestNumber = (int) number;
            }
            else
            {
                argList[ argCount ] = arg;
                argCount++;
            }
        }
    }

    // Check for the wimpy mode environment variable
    if (getenv("CL_WIMPY_MODE")) {
      vlog( "\n" );
      vlog( "*** Detected CL_WIMPY_MODE env                          ***\n" );
      gWimpyMode = 1;
    }

    vlog( "\n" );

    vlog( "Test binary built %s %s\n", __DATE__, __TIME__ );

    PrintArch();

    if( gWimpyMode )
    {
        vlog( "\n" );
        vlog( "*** WARNING: Testing in Wimpy mode!                     ***\n" );
        vlog( "*** Wimpy mode is not sufficient to verify correctness. ***\n" );
        vlog( "*** It gives warm fuzzy feelings and then nevers calls. ***\n\n" );
        vlog("*** Wimpy Reduction Factor: %-27u ***\n\n", gWimpyReductionFactor);
    }

    return 0;
}

static void PrintUsage( void )
{
    int i;
    vlog( "%s [-wz#]: <optional: test names>\n", appName );
    vlog( "\ttest names:\n" );
    vlog( "\t\tdestFormat<_sat><_round>_sourceFormat\n" );
    vlog( "\t\t\tPossible format types are:\n\t\t\t\t" );
    for( i = 0; i < kTypeCount; i++ )
        vlog( "%s, ", gTypeNames[i] );
    vlog( "\n\n\t\t\tPossible saturation values are: (empty) and _sat\n" );
    vlog( "\t\t\tPossible rounding values are:\n\t\t\t\t(empty), " );
    for( i = 1; i < kRoundingModeCount; i++ )
        vlog( "%s, ", gRoundingModeNames[i] );
    vlog( "\n\t\t\tExamples:\n" );
    vlog( "\t\t\t\tulong_short   converts short to ulong\n" );
    vlog( "\t\t\t\tchar_sat_rte_float   converts float to char with saturated clipping in round to nearest rounding mode\n\n" );
    vlog( "\toptions:\n" );
    vlog( "\t\t-d\tToggle testing of double precision.  On by default if cl_khr_fp64 is enabled, ignored otherwise.\n" );
    vlog( "\t\t-l\tToggle link check mode. When on, testing is skipped, and we just check to see that the kernels build. (Off by default.)\n" );
    vlog( "\t\t-m\tToggle Multithreading. (On by default.)\n" );
    vlog( "\t\t-w\tToggle wimpy mode. When wimpy mode is on, we run a very small subset of the tests for each fn. NOT A VALID TEST! (Off by default.)\n" );
    vlog(" \t\t-[2^n]\tSet wimpy reduction factor, recommended range of n is 1-12, default factor(%u)\n", gWimpyReductionFactor);
    vlog( "\t\t-z\tToggle flush to zero mode  (Default: per device)\n" );
    vlog( "\t\t-#\tTest just vector size given by #, where # is an element of the set {1,2,3,4,8,16}\n" );
    vlog( "\n" );
    vlog( "You may also pass the number of the test on which to start.\nA second number can be then passed to indicate how many tests to run\n\n" );
}


static int GetTestCase( const char *name, Type *outType, Type *inType, SaturationMode *sat, RoundingMode *round )
{
    int i;

    //Find the return type
    for( i = 0; i < kTypeCount; i++ )
        if( name == strstr( name, gTypeNames[i] ) )
        {
            *outType = (Type)i;
            name += strlen( gTypeNames[i] );

            break;
        }

    if( i == kTypeCount )
        return -1;

    // Check to see if _sat appears next
    *sat = (SaturationMode)0;
    for( i = 1; i < kSaturationModeCount; i++ )
        if( name == strstr( name, gSaturationNames[i] ) )
        {
            *sat = (SaturationMode)i;
            name += strlen( gSaturationNames[i] );
            break;
        }

    *round = (RoundingMode)0;
    for( i = 1; i < kRoundingModeCount; i++ )
        if( name == strstr( name, gRoundingModeNames[i] ) )
        {
            *round = (RoundingMode)i;
            name += strlen( gRoundingModeNames[i] );
            break;
        }

    if( *name != '_' )
        return -2;
    name++;

    for( i = 0; i < kTypeCount; i++ )
        if( name == strstr( name, gTypeNames[i] ) )
        {
            *inType = (Type)i;
            name += strlen( gTypeNames[i] );

            break;
        }

    if( i == kTypeCount )
        return -3;

    if( *name != '\0' )
        return -4;

    return 0;
}

#pragma mark -
#pragma mark OpenCL

test_status InitCL( cl_device_id device )
{
    int error, i;
    size_t configSize = sizeof( gComputeDevices );

    if( (error = clGetDeviceInfo( device, CL_DEVICE_MAX_COMPUTE_UNITS, configSize, &gComputeDevices, NULL )) )
        gComputeDevices = 1;

    configSize = sizeof( gDeviceFrequency );
    if( (error = clGetDeviceInfo( device, CL_DEVICE_MAX_CLOCK_FREQUENCY, configSize, &gDeviceFrequency, NULL )) )
        gDeviceFrequency = 0;

    cl_device_fp_config floatCapabilities = 0;
    if( (error = clGetDeviceInfo(device, CL_DEVICE_SINGLE_FP_CONFIG, sizeof(floatCapabilities), &floatCapabilities,  NULL)))
        floatCapabilities = 0;
    if(0 == (CL_FP_DENORM & floatCapabilities) )
        gForceFTZ ^= 1;

    if( 0 == (floatCapabilities & CL_FP_ROUND_TO_NEAREST ) )
    {
        char profileStr[128] = "";
        // Verify that we are an embedded profile device
        if( (error = clGetDeviceInfo( device, CL_DEVICE_PROFILE, sizeof( profileStr ), profileStr, NULL ) ) )
        {
            vlog_error( "FAILURE: Could not get device profile: error %d\n", error );
            return TEST_FAIL;
        }

        if( strcmp( profileStr, "EMBEDDED_PROFILE" ) )
        {
            vlog_error( "FAILURE: non-embedded profile device does not support CL_FP_ROUND_TO_NEAREST\n" );
            return TEST_FAIL;
        }

        if( 0 == (floatCapabilities & CL_FP_ROUND_TO_ZERO ) )
        {
            vlog_error( "FAILURE: embedded profile device supports neither CL_FP_ROUND_TO_NEAREST or CL_FP_ROUND_TO_ZERO\n" );
            return TEST_FAIL;
        }

        gIsRTZ = 1;
    }

    else if(is_extension_available(device, "cl_khr_fp64"))
    {
        gHasDouble = 1;
    }
    gTestDouble &= gHasDouble;

    //detect whether profile of the device is embedded
    char profile[1024] = "";
    if( (error = clGetDeviceInfo( device, CL_DEVICE_PROFILE, sizeof(profile), profile, NULL ) ) ){}
    else if( strstr(profile, "EMBEDDED_PROFILE" ) )
    {
        gIsEmbedded = 1;
        if( !is_extension_available(device, "cles_khr_int64" ) )
            gHasLong = 0;
    }


    gContext = clCreateContext( NULL, 1, &device, notify_callback, NULL, &error );
    if( NULL == gContext || error )
    {
        vlog_error( "clCreateContext failed. (%d)\n", error );
        return TEST_FAIL;
    }

    gQueue = clCreateCommandQueue(gContext, device, 0, &error);
    if( NULL == gQueue || error )
    {
        vlog_error( "clCreateCommandQueue failed. (%d)\n", error );
        return TEST_FAIL;
    }

    //Allocate buffers
    //FIXME: use clProtectedArray for guarded allocations?
    gIn   = malloc( BUFFER_SIZE + 2 * kPageSize );
    gAllowZ = malloc( BUFFER_SIZE + 2 * kPageSize );
    gRef  = malloc( BUFFER_SIZE + 2 * kPageSize );
    for( i = 0; i < kCallStyleCount; i++ )
    {
        gOut[i] = malloc( BUFFER_SIZE + 2 * kPageSize );
        if( NULL == gOut[i] )
            return TEST_FAIL;
    }

    // setup input buffers
    gInBuffer = clCreateBuffer(gContext, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, BUFFER_SIZE, NULL, &error);
    if( gInBuffer == NULL || error)
    {
        vlog_error( "clCreateBuffer failed for input (%d)\n", error );
        return TEST_FAIL;
    }

    // setup output buffers
    for( i = 0; i < kCallStyleCount; i++ )
    {
        gOutBuffers[i] = clCreateBuffer(  gContext, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, BUFFER_SIZE, NULL, &error );
        if( gOutBuffers[i] == NULL || error )
        {
            vlog_error( "clCreateArray failed for output (%d)\n", error );
            return TEST_FAIL;
        }
    }


    gMTdata = init_genrand( gRandomSeed );


    char c[1024];
    static const char *no_yes[] = { "NO", "YES" };
    vlog( "\nCompute Device info:\n" );
    clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(c), c, NULL);
    vlog( "\tDevice Name: %s\n", c );
    clGetDeviceInfo(device, CL_DEVICE_VENDOR, sizeof(c), c, NULL);
    vlog( "\tVendor: %s\n", c );
    clGetDeviceInfo(device, CL_DEVICE_VERSION, sizeof(c), c, NULL);
    vlog( "\tDevice Version: %s\n", c );
    clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_VERSION, sizeof(c), &c, NULL);
    vlog( "\tCL C Version: %s\n", c );
    clGetDeviceInfo(device, CL_DRIVER_VERSION, sizeof(c), c, NULL);
    vlog( "\tDriver Version: %s\n", c );
    vlog( "\tProcessing with %ld devices\n", gComputeDevices );
    vlog( "\tDevice Frequency: %d MHz\n", gDeviceFrequency );
    vlog( "\tSubnormal values supported for floats? %s\n", no_yes[0 != (CL_FP_DENORM & floatCapabilities)] );
    vlog( "\tTesting with FTZ mode ON for floats? %s\n", no_yes[0 != gForceFTZ] );
    vlog( "\tTesting with default RTZ mode for floats? %s\n", no_yes[0 != gIsRTZ] );
    vlog( "\tHas Double? %s\n", no_yes[0 != gHasDouble] );
    if( gHasDouble )
        vlog( "\tTest Double? %s\n", no_yes[0 != gTestDouble] );
    vlog( "\tHas Long? %s\n", no_yes[0 != gHasLong] );
    vlog( "\tTesting vector sizes: " );
    for( i = gMinVectorSize; i < gMaxVectorSize; i++ )
        vlog("\t%d", vectorSizes[i]);
    vlog( "\n" );
    return TEST_PASS;
}

static int RunKernel( cl_kernel kernel, void *inBuf, void *outBuf, size_t blockCount )
{
    // The global dimensions are just the blockCount to execute since we haven't set up multiple queues for multiple devices.
    int error;

    error = clSetKernelArg(kernel, 0, sizeof( inBuf ), &inBuf);
    error |= clSetKernelArg(kernel, 1, sizeof(outBuf), &outBuf);

    if( error )
    {
        vlog_error( "FAILED -- could not set kernel args (%d)\n", error );
        return error;
    }

    if( (error = clEnqueueNDRangeKernel(gQueue, kernel, 1, NULL, &blockCount, NULL, 0, NULL, NULL)))
    {
        vlog_error( "FAILED -- could not execute kernel (%d)\n", error );
        return error;
    }

    return 0;
}

#if ! defined( __APPLE__ )
void memset_pattern4(void *dest, const void *src_pattern, size_t bytes );
#endif

#if defined( __APPLE__ )
#include <mach/mach_time.h>
#endif

uint64_t GetTime( void );
uint64_t GetTime( void )
{
#if defined( __APPLE__ )
    return mach_absolute_time();
#elif defined(_MSC_VER)
    return  ReadTime();
#else
    //mach_absolute_time is a high precision timer with precision < 1 microsecond.
#warning need accurate clock here.  Times are invalid.
    return 0;
#endif
}


#if defined (_MSC_VER)
/* function is defined in "compat.h" */
#else
double SubtractTime( uint64_t endTime, uint64_t startTime );
double SubtractTime( uint64_t endTime, uint64_t startTime )
{
    uint64_t diff = endTime - startTime;
    static double conversion = 0.0;

    if( 0.0 == conversion )
    {
#if defined( __APPLE__ )
        mach_timebase_info_data_t info = {0,0};
        kern_return_t   err = mach_timebase_info( &info );
        if( 0 == err )
            conversion = 1e-9 * (double) info.numer / (double) info.denom;
#else
        // This function consumes output from GetTime() above, and converts the time to secionds.
#warning need accurate ticks to seconds conversion factor here. Times are invalid.
#endif
    }

    // strictly speaking we should also be subtracting out timer latency here
    return conversion * (double) diff;
}
#endif

typedef struct CalcReferenceValuesInfo
{
    struct WriteInputBufferInfo *parent;        // pointer back to the parent WriteInputBufferInfo struct
    cl_kernel                   kernel;         // the kernel for this vector size
    cl_program                  program;        // the program for this vector size
    cl_uint                     vectorSize;     // the vector size for this callback chain
    void                        *p;             // the pointer to mapped result data for this vector size
    cl_int                      result;
}CalcReferenceValuesInfo;

typedef struct WriteInputBufferInfo
{
    volatile cl_event           calcReferenceValues;   // user event which signals when main thread is done calculating reference values
    volatile cl_event           doneBarrier;     // user event which signals when worker threads are done
    cl_uint                     count;           // the number of elements in the array
    Type                        outType;         // the data type of the conversion result
    Type                        inType;          // the data type of the conversion input
    volatile int                barrierCount;
    CalcReferenceValuesInfo     calcInfo[kCallStyleCount];
}WriteInputBufferInfo;

cl_uint RoundUpToNextPowerOfTwo( cl_uint x );
cl_uint RoundUpToNextPowerOfTwo( cl_uint x )
{
    if( 0 == (x & (x-1)))
        return x;

    while( x & (x-1) )
       x &= x-1;

    return x + x;
}

void CL_CALLBACK WriteInputBufferComplete( cl_event, cl_int, void * );

typedef struct DataInitInfo
{
    cl_ulong        start;
    cl_uint         size;
    Type            outType;
    Type            inType;
    SaturationMode  sat;
    RoundingMode    round;
    MTdata          *d;
}DataInitInfo;

cl_int InitData( cl_uint job_id, cl_uint thread_id, void *p );
cl_int InitData( cl_uint job_id, cl_uint thread_id, void *p )
{
    DataInitInfo *info = (DataInitInfo*) p;

    gInitFunctions[ info->inType ]( (char*)gIn + job_id * info->size * gTypeSizes[info->inType], info->sat, info->round,
                                   info->outType, info->start + job_id * info->size, info->size, info->d[thread_id] );
    return CL_SUCCESS;
}

static void setAllowZ(uint8_t *allow, uint32_t *x, cl_uint count)
{
    cl_uint i;
    for (i = 0; i < count; ++i)
    allow[i] |= (uint8_t)((x[i] & 0x7f800000U) == 0);
}

cl_int PrepareReference( cl_uint job_id, cl_uint thread_id, void *p );
cl_int PrepareReference( cl_uint job_id, cl_uint thread_id, void *p )
{
    DataInitInfo *info = (DataInitInfo*) p;
    cl_uint count = info->size;
    Type inType = info->inType;
    Type outType = info->outType;
    RoundingMode round = info->round;
    size_t j;

    Force64BitFPUPrecision();

    void *s = (cl_uchar*) gIn + job_id * count * gTypeSizes[info->inType];
    void *a = (cl_uchar*) gAllowZ + job_id * count;
    void *d = (cl_uchar*) gRef + job_id * count * gTypeSizes[info->outType];

    if (outType != inType)
    {
        //create the reference while we wait
        Convert f = gConversions[ outType ][ inType ];
        if( info->sat )
            f = gSaturatedConversions[ outType ][ inType ];

#if (defined(__arm__) || defined(__aarch64__)) && defined(__GNUC__)
        /* ARM VFP doesn't have hardware instruction for converting from 64-bit
         * integer to float types, hence GCC ARM uses the floating-point
         * emulation code despite which -mfloat-abi setting it is. But the
         * emulation code in libgcc.a has only one rounding mode (round to
         * nearest even in this case) and ignores the user rounding mode setting
         * in hardware. As a result setting rounding modes in hardware won't
         * give correct rounding results for type covert from 64-bit integer to
         * float using GCC for ARM compiler so for testing different rounding
         * modes, we need to use alternative reference function. ARM64 does have
         * an instruction, however we cannot guarantee the compiler will use it.
         * On all ARM architechures use emulation to calculate reference.*/
        switch (round)
        {
            /* conversions to floating-point type use the current rounding mode.
             * The only default floating-point rounding mode supported is round to nearest even
             * i.e the current rounding mode will be _rte for floating-point types. */
            case kDefaultRoundingMode:
                    qcom_rm = qcomRTE;
                    break;
            case kRoundToNearestEven:
                    qcom_rm = qcomRTE;
                    break;
            case kRoundUp:
                    qcom_rm = qcomRTP;
                    break;
            case kRoundDown:
                    qcom_rm = qcomRTN;
                    break;
            case kRoundTowardZero:
                    qcom_rm = qcomRTZ;
                    break;
            default:
                    vlog_error("ERROR: undefined rounding mode %d\n", round);
                    break;
        }
        qcom_sat =  info->sat;
#endif

        RoundingMode oldRound = set_round( round, outType );
        f( d, s, count );
        set_round( oldRound, outType );

    // Decide if we allow a zero result in addition to the correctly rounded one
        memset(a, 0, count);
    if (gForceFTZ) {
        if (inType == kfloat)
        setAllowZ((uint8_t*)a, (uint32_t*)s, count);
        if (outType == kfloat)
        setAllowZ((uint8_t*)a, (uint32_t*)d, count);
    }
    }
    else
    {
        // Copy the input to the reference
        memcpy(d, s, info->size * gTypeSizes[inType]);
    }

    //Patch up NaNs conversions to integer to zero -- these can be converted to any integer
    if( info->outType != kfloat && info->outType != kdouble )
    {
        if( inType == kfloat )
        {
            float *inp = (float*) s;
            for( j = 0; j < count; j++ )
            {
                if( isnan( inp[j] ) )
                    memset( (char*) d + j * gTypeSizes[ outType ], 0, gTypeSizes[ outType ] );
            }
        }
        if( inType == kdouble )
        {
            double *inp = (double*) s;
            for( j = 0; j < count; j++ )
            {
                if( isnan( inp[j] ) )
                    memset( (char*) d + j * gTypeSizes[ outType ], 0, gTypeSizes[ outType ] );
            }
        }
    }
    else if( inType == kfloat || inType == kdouble )
    {  // outtype and intype is float or double.  NaN conversions for float <-> double can be any NaN
        if( inType == kfloat && outType == kdouble )
        {
            float *inp = (float*) s;
            for( j = 0; j < count; j++ )
            {
                if( isnan( inp[j] ) )
                    ((double*) d)[j] = NAN;
            }
        }
        if( inType == kdouble && outType == kfloat )
        {
            double *inp = (double*) s;
            for( j = 0; j < count; j++ )
            {
                if( isnan( inp[j] ) )
                    ((float*) d)[j] = NAN;
            }
        }
    }

    return CL_SUCCESS;
}

static int DoTest( cl_device_id device, Type outType, Type inType, SaturationMode sat, RoundingMode round, MTdata d )
{
#ifdef __APPLE__
    cl_ulong wall_start = mach_absolute_time();
#endif

    DataInitInfo  init_info = { 0, 0, outType, inType, sat, round, NULL };
    WriteInputBufferInfo writeInputBufferInfo;
    int vectorSize;
    int error = 0;
    cl_uint threads = GetThreadCount();
    uint64_t i;

    gTestCount++;
    size_t blockCount = BUFFER_SIZE / MAX( gTypeSizes[ inType ], gTypeSizes[ outType ] );
    size_t step = blockCount;
    uint64_t lastCase = 1ULL << (8*gTypeSizes[ inType ]);
    cl_event writeInputBuffer = NULL;

    memset( &writeInputBufferInfo, 0, sizeof( writeInputBufferInfo ) );
    init_info.d = (MTdata*)malloc( threads * sizeof( MTdata ) );
    if( NULL == init_info.d )
    {
        vlog_error( "ERROR: Unable to allocate storage for random number generator!\n" );
        return -1;
    }
    for( i = 0; i < threads; i++ )
    {
        init_info.d[i] = init_genrand( genrand_int32( d ) );
        if( NULL == init_info.d[i] )
        {
            vlog_error( "ERROR: Unable to allocate storage for random number generator!\n" );
            return -1;
        }
    }

    writeInputBufferInfo.outType = outType;
    writeInputBufferInfo.inType = inType;

    for( vectorSize = gMinVectorSize; vectorSize < gMaxVectorSize; vectorSize++)
    {
        writeInputBufferInfo.calcInfo[vectorSize].program = MakeProgram( outType, inType, sat, round, vectorSize,
                                                                        &writeInputBufferInfo.calcInfo[vectorSize].kernel );
        if( NULL == writeInputBufferInfo.calcInfo[vectorSize].program )
        {
            gFailCount++;
            return -1;
        }
        if( NULL == writeInputBufferInfo.calcInfo[vectorSize].kernel )
        {
            gFailCount++;
            vlog_error( "\t\tFAILED -- Failed to create kernel.\n" );
            return -2;
        }

        writeInputBufferInfo.calcInfo[vectorSize].parent = &writeInputBufferInfo;
        writeInputBufferInfo.calcInfo[vectorSize].vectorSize = vectorSize;
        writeInputBufferInfo.calcInfo[vectorSize].result = -1;
    }

    if( gSkipTesting )
        goto exit;

    // Patch up rounding mode if default is RTZ
    // We leave the part above in default rounding mode so that the right kernel is compiled.
    if( round == kDefaultRoundingMode && gIsRTZ && (outType == kfloat) )
        init_info.round = round = kRoundTowardZero;

    // Figure out how many elements are in a work block

    // we handle 64-bit types a bit differently.
    if( 8*gTypeSizes[ inType ] > 32 )
        lastCase = 0x100000000ULL;

    if ( !gWimpyMode && gIsEmbedded )
      step = blockCount * EMBEDDED_REDUCTION_FACTOR;

    if ( gWimpyMode )
        step = (size_t)blockCount * (size_t)gWimpyReductionFactor;
    vlog( "Testing... " );
    fflush(stdout);
    for( i = 0; i < (uint64_t)lastCase; i += step )
    {

        if( 0 == ( i & ((lastCase >> 3) -1))) {
            vlog(".");
            fflush(stdout);
        }

        cl_uint count = (uint32_t) MIN( blockCount, lastCase - i );
        writeInputBufferInfo.count = count;

        // Crate a user event to represent the status of the reference value computation completion
        writeInputBufferInfo.calcReferenceValues = clCreateUserEvent( gContext, &error);
        if( error || NULL == writeInputBufferInfo.calcReferenceValues )
        {
            vlog_error( "ERROR: Unable to create user event. (%d)\n", error );
            gFailCount++;
            goto exit;
        }

        // retain for consumption by MapOutputBufferComplete
        for( vectorSize = gMinVectorSize; vectorSize < gMaxVectorSize; vectorSize++)
        {
            if( (error = clRetainEvent(writeInputBufferInfo.calcReferenceValues) ))
            {
                vlog_error( "ERROR: Unable to retain user event. (%d)\n", error );
                gFailCount++;
                goto exit;
            }
        }

        // Crate a user event to represent when the callbacks are done verifying correctness
        writeInputBufferInfo.doneBarrier = clCreateUserEvent( gContext, &error);
        if( error || NULL == writeInputBufferInfo.calcReferenceValues )
        {
            vlog_error( "ERROR: Unable to create user event for barrier. (%d)\n", error );
            gFailCount++;
            goto exit;
        }

        // retain for use by the callback that calls this
        if( (error = clRetainEvent(writeInputBufferInfo.doneBarrier) ))
        {
            vlog_error( "ERROR: Unable to retain user event doneBarrier. (%d)\n", error );
            gFailCount++;
            goto exit;
        }

        //      Call this in a multithreaded manner
        //      gInitFunctions[ inType ]( gIn, sat, round, outType, i, count, d );
        cl_uint chunks = RoundUpToNextPowerOfTwo(threads) * 2;
        init_info.start = i;
        init_info.size = count / chunks;
        if( init_info.size < 16384 )
        {
            chunks = RoundUpToNextPowerOfTwo(threads);
            init_info.size = count / chunks;
            if( init_info.size < 16384 )
            {
                init_info.size = count;
                chunks = 1;
            }
        }
        ThreadPool_Do(InitData, chunks, &init_info);

        // Copy the results to the device
        writeInputBuffer = NULL;
        if( (error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0, count * gTypeSizes[inType], gIn, 0, NULL, &writeInputBuffer )))
        {
            vlog_error( "ERROR: clEnqueueWriteBuffer failed. (%d)\n", error );
            gFailCount++;
            goto exit;
        }

        // Setup completion callback for the write, which will enqueue the rest of the work
        // This is somewhat gratuitous.  Because this is an in order queue, we didn't really need to
        // do this work in a callback. We could have done it from the main thread.  Here we are
        // verifying that the implementation can enqueue work from a callback, while at the same time
        // also checking to make sure that the conversions work.
        //
        // Because the verification code is also moved to a callback, it is hoped that implementations will
        // achieve a test performance improvement because they can verify the results in parallel.  If the
        // implementation serializes callbacks however, that won't happen.   Consider it some motivation
        // to do the right thing! :-)
        if( (error = clSetEventCallback( writeInputBuffer, CL_COMPLETE, WriteInputBufferComplete, &writeInputBufferInfo)) )
        {
            vlog_error( "ERROR: clSetEventCallback failed. (%d)\n", error );
            gFailCount++;
            goto exit;
        }

        // The event can't be destroyed until the callback is called, so we can release it now.
        if( (error = clReleaseEvent(writeInputBuffer) ))
        {
            vlog_error( "ERROR: clReleaseEvent failed. (%d)\n", error );
            gFailCount++;
            goto exit;
        }

        // Make sure the work is actually running, so we don't deadlock
        if( (error = clFlush( gQueue ) ) )
        {
            vlog_error( "clFlush failed with error %d\n", error );
            gFailCount++;
            goto exit;
        }

        ThreadPool_Do(PrepareReference, chunks, &init_info);

        // signal we are done calculating the reference results
        if( (error = clSetUserEventStatus( writeInputBufferInfo.calcReferenceValues, CL_COMPLETE ) ) )
        {
            vlog_error( "Error:  Failed to set user event status to CL_COMPLETE:  %d\n", error );
            gFailCount++;
            goto exit;
        }

        // Wait for the event callbacks to finish verifying correctness.
        if( (error = clWaitForEvents( 1, (cl_event*) &writeInputBufferInfo.doneBarrier ) ))
        {
            vlog_error( "Error:  Failed to wait for barrier:  %d\n", error );
            gFailCount++;
            goto exit;
        }

        if( (error = clReleaseEvent(writeInputBufferInfo.calcReferenceValues ) ))
        {
            vlog_error( "Error:  Failed to release calcReferenceValues:  %d\n", error );
            gFailCount++;
            goto exit;
        }

        if( (error = clReleaseEvent(writeInputBufferInfo.doneBarrier ) ))
        {
            vlog_error( "Error:  Failed to release done barrier:  %d\n", error );
            gFailCount++;
            goto exit;
        }


        for( vectorSize = gMinVectorSize; vectorSize < gMaxVectorSize; vectorSize++)
        {
            if( ( error = writeInputBufferInfo.calcInfo[ vectorSize ].result ))
            {
                switch( inType )
                {
                    case kuchar:
                    case kchar:
                        vlog( "Input value: 0x%2.2x ", ((unsigned char*)gIn)[error - 1] );
                        break;
                    case kushort:
                    case kshort:
                        vlog( "Input value: 0x%4.4x ", ((unsigned short*)gIn)[error - 1] );
                        break;
                    case kuint:
                    case kint:
                        vlog( "Input value: 0x%8.8x ", ((unsigned int*)gIn)[error - 1] );
                        break;
                    case kfloat:
                        vlog( "Input value: %a ", ((float*)gIn)[error - 1] );
                        break;
                        break;
                    case kulong:
                    case klong:
                        vlog( "Input value: 0x%16.16llx ", ((unsigned long long*)gIn)[error - 1] );
                        break;
                    case kdouble:
                        vlog( "Input value: %a ", ((double*)gIn)[error - 1]);
                        break;
                    default:
                        vlog_error( "Internal error at %s: %d\n", __FILE__, __LINE__ );
                        abort();
                        break;
                }

                // tell the user which conversion it was.
                if( 0 == vectorSize )
                    vlog( " (implicit scalar conversion from %s to %s)\n", gTypeNames[ inType ], gTypeNames[ outType ] );
                else
                    vlog( " (convert_%s%s%s%s( %s%s ))\n", gTypeNames[outType], sizeNames[vectorSize], gSaturationNames[ sat ],
                                                            gRoundingModeNames[ round ], gTypeNames[inType], sizeNames[vectorSize] );

                gFailCount++;
                goto exit;
            }
        }
    }

    log_info( "done.\n" );

    if( gTimeResults )
    {
        //Kick off tests for the various vector lengths
        for( vectorSize = gMinVectorSize; vectorSize < gMaxVectorSize; vectorSize++)
        {
            size_t workItemCount = blockCount / vectorSizes[vectorSize];
            if( vectorSizes[vectorSize] * gTypeSizes[outType] < 4 )
                workItemCount /= 4 / (vectorSizes[vectorSize] * gTypeSizes[outType]);

            double sum = 0.0;
            double bestTime = INFINITY;
            cl_uint k;
            for( k = 0; k < PERF_LOOP_COUNT; k++ )
            {
                uint64_t startTime = GetTime();
                if( (error = RunKernel( writeInputBufferInfo.calcInfo[vectorSize].kernel, gInBuffer, gOutBuffers[ vectorSize ], workItemCount )) )
                {
                    gFailCount++;
                    goto exit;
                }

                // Make sure OpenCL is done
                if( (error = clFinish(gQueue) ) )
                {
                    vlog_error( "Error %d at clFinish\n", error );
                    goto exit;
                }

                uint64_t endTime = GetTime();
                double time = SubtractTime( endTime, startTime );
                sum += time;
                if( time < bestTime )
                    bestTime = time;

            }

            if( gReportAverageTimes )
                bestTime = sum / PERF_LOOP_COUNT;
            double clocksPerOp = bestTime * (double) gDeviceFrequency * gComputeDevices * gSimdSize * 1e6 / (workItemCount * vectorSizes[vectorSize]);
            if( 0 == vectorSize )
                vlog_perf( clocksPerOp, LOWER_IS_BETTER, "clocks / element", "implicit convert %s -> %s", gTypeNames[ inType ], gTypeNames[ outType ] );
            else
                vlog_perf( clocksPerOp, LOWER_IS_BETTER, "clocks / element", "convert_%s%s%s%s( %s%s )", gTypeNames[ outType ], sizeNames[vectorSize], gSaturationNames[ sat ], gRoundingModeNames[round], gTypeNames[inType], sizeNames[vectorSize] );
        }
    }

    if( gWimpyMode )
        vlog( "\tWimp pass" );
    else
        vlog( "\tpassed" );

#ifdef __APPLE__
    // record the run time
    vlog( "\t(%f s)", 1e-9 * ( mach_absolute_time() - wall_start ) );
#endif
    vlog( "\n\n" );
    fflush( stdout );


exit:
    //clean up
    for( vectorSize = gMinVectorSize; vectorSize < gMaxVectorSize; vectorSize++)
    {
        clReleaseProgram( writeInputBufferInfo.calcInfo[vectorSize].program );
        clReleaseKernel( writeInputBufferInfo.calcInfo[vectorSize].kernel );
    }

    if( init_info.d )
    {
        for( i = 0; i < threads; i++ )
            free_mtdata(init_info.d[i]);
        free(init_info.d);
    }

    return error;
}

void CL_CALLBACK MapResultValuesComplete( cl_event e, cl_int status, void *data );

// Note: not called reentrantly
void CL_CALLBACK WriteInputBufferComplete( cl_event e, cl_int status, void *data )
{
    WriteInputBufferInfo *info = (WriteInputBufferInfo*) data;
    cl_uint count = info->count;
    int vectorSize;

    if( CL_SUCCESS != status )
    {
        vlog_error( "ERROR: WriteInputBufferComplete calback failed with status: %d\n", status );
        gFailCount++;
        return;
    }

    info->barrierCount = gMaxVectorSize - gMinVectorSize;

    // now that we know that the write buffer is complete, enqueue callbacks to wait for the main thread to
    // finish calculating the reference results.
    for( vectorSize = gMinVectorSize; vectorSize < gMaxVectorSize; vectorSize++)
    {
        size_t workItemCount = (count + vectorSizes[vectorSize] - 1) / ( vectorSizes[vectorSize]);
        cl_event mapComplete = NULL;

        if( (status = RunKernel( info->calcInfo[ vectorSize ].kernel, gInBuffer, gOutBuffers[ vectorSize ], workItemCount )) )
        {
            gFailCount++;
            return;
        }

        info->calcInfo[vectorSize].p = clEnqueueMapBuffer( gQueue, gOutBuffers[ vectorSize ], CL_FALSE, CL_MAP_READ | CL_MAP_WRITE,
                                                          0, count * gTypeSizes[ info->outType ], 0, NULL, &mapComplete, &status);
        {
            if( status )
            {
                vlog_error( "ERROR: WriteInputBufferComplete calback failed with status: %d\n", status );
                gFailCount++;
                return;
            }
        }

        if( (status = clSetEventCallback( mapComplete, CL_COMPLETE, MapResultValuesComplete, info->calcInfo + vectorSize)))
        {
            vlog_error( "ERROR: WriteInputBufferComplete calback failed with status: %d\n", status );
            gFailCount++;
            return;
        }

        if( (status = clReleaseEvent(mapComplete)))
        {
            vlog_error( "ERROR: clReleaseEvent calback failed in WriteInputBufferComplete for vector size %d with status: %d\n", vectorSize, status );
            gFailCount++;
            return;
        }
    }

    // Make sure the work starts moving -- otherwise we may deadlock
    if( (status = clFlush(gQueue)))
    {
        vlog_error( "ERROR: WriteInputBufferComplete calback failed with status: %d\n", status );
        gFailCount++;
        return;
    }

    // e was already released by the main thread. It should be destroyed automatically soon after we exit.
}

void CL_CALLBACK CalcReferenceValuesComplete( cl_event e, cl_int status, void *data );

// Note: May be called reentrantly
void CL_CALLBACK MapResultValuesComplete( cl_event e, cl_int status, void *data )
{
    CalcReferenceValuesInfo *info = (CalcReferenceValuesInfo*) data;
    cl_event calcReferenceValues = info->parent->calcReferenceValues;

    if( CL_SUCCESS != status )
    {
        vlog_error( "ERROR: MapResultValuesComplete calback failed with status: %d\n", status );
        gFailCount++;       // not thread safe -- being lazy here
        clReleaseEvent(calcReferenceValues);
        return;
    }

    // we know that the map is done, wait for the main thread to finish calculating the reference values
    if( (status = clSetEventCallback( calcReferenceValues, CL_COMPLETE, CalcReferenceValuesComplete, data )))
    {
        vlog_error( "ERROR: clSetEventCallback failed in MapResultValuesComplete with status: %d\n", status );
        gFailCount++;       // not thread safe -- being lazy here
    }

    // this thread no longer needs its reference to info->calcReferenceValues, so release it
    if( (status = clReleaseEvent(calcReferenceValues) ))
    {
        vlog_error( "ERROR: clReleaseEvent(info->calcReferenceValues) failed with status: %d\n", status );
        gFailCount++;       // not thread safe -- being lazy here
    }

    // no need to flush since we didn't enqueue anything

    // e was already released by WriteInputBufferComplete. It should be destroyed automatically soon after we exit.
}


void CL_CALLBACK CalcReferenceValuesComplete( cl_event e, cl_int status, void *data )
{
    CalcReferenceValuesInfo     *info = (CalcReferenceValuesInfo*) data;
    cl_uint                     vectorSize = info->vectorSize;
    cl_uint                     count = info->parent->count;
    Type                        outType = info->parent->outType;        // the data type of the conversion result
    Type                        inType = info->parent->inType;          // the data type of the conversion input
    size_t                      j;
    cl_int                      error;
    cl_event                    doneBarrier = info->parent->doneBarrier;

    // report spurious error condition
    if( CL_SUCCESS != status )
    {
        vlog_error( "ERROR: CalcReferenceValuesComplete did not succeed! (%d)\n", status );
        gFailCount++;       // lazy about thread safety here
        return;
    }

    // Now we know that both results have been mapped back from the device, and the
    // main thread is done calculating the reference results. It is now time to check
    // the results.

    // verify results
    void *mapped = info->p;

    //Patch up NaNs conversions to integer to zero -- these can be converted to any integer
    if( outType != kfloat && outType != kdouble )
    {
        if( inType == kfloat )
        {
            float *inp = (float*) gIn;
            for( j = 0; j < count; j++ )
            {
                if( isnan( inp[j] ) )
                    memset( (char*) mapped + j * gTypeSizes[ outType ], 0, gTypeSizes[ outType ] );
            }
        }
        if( inType == kdouble )
        {
            double *inp = (double*) gIn;
            for( j = 0; j < count; j++ )
            {
                if( isnan( inp[j] ) )
                    memset( (char*) mapped + j * gTypeSizes[ outType ], 0, gTypeSizes[ outType ] );
            }
        }
    }
    else if( inType == kfloat || inType == kdouble )
    {  // outtype and intype is float or double.  NaN conversions for float <-> double can be any NaN
        if( inType == kfloat && outType == kdouble )
        {
            float *inp = (float*) gIn;
            double *outp = (double*) mapped;
            for( j = 0; j < count; j++ )
            {
                if( isnan( inp[j] ) && isnan(outp[j]) )
                    outp[j] = NAN;
            }
        }
        if( inType == kdouble && outType == kfloat )
        {
            double *inp = (double*) gIn;
            float *outp = (float*) mapped;
            for( j = 0; j < count; j++ )
            {
                if( isnan( inp[j] ) && isnan(outp[j]) )
                    outp[j] = NAN;
            }
        }
    }

    if( memcmp( mapped, gRef, count * gTypeSizes[ outType ] ) )
        info->result = gCheckResults[outType]( mapped, gRef, gAllowZ, count, vectorSizes[vectorSize] );
    else
        info->result = 0;

    // Fill the output buffer with junk and release it
    {
        cl_uint pattern =  0xffffdead;
        memset_pattern4(mapped, &pattern, count * gTypeSizes[outType]);
        if((error = clEnqueueUnmapMemObject(gQueue, gOutBuffers[ vectorSize ], mapped, 0, NULL, NULL)))
        {
            vlog_error( "ERROR: clEnqueueUnmapMemObject failed in CalcReferenceValuesComplete  (%d)\n", error );
            gFailCount++;
        }
    }

    if( 1 == ThreadPool_AtomicAdd( &info->parent->barrierCount, -1) )
    {
        if( (status = clSetUserEventStatus( doneBarrier, CL_COMPLETE) ))
        {
            vlog_error( "ERROR: clSetUserEventStatus failed in CalcReferenceValuesComplete (err: %d). We're probably going to deadlock.\n", status );
            gFailCount++;
            return;
        }

        if( (status = clReleaseEvent( doneBarrier ) ) )
        {
            vlog_error( "ERROR: clReleaseEvent failed in CalcReferenceValuesComplete (err: %d).\n", status );
            gFailCount++;
            return;
        }
    }


    // e was already released by WriteInputBufferComplete. It should be destroyed automatically soon after
    // all the calls to CalcReferenceValuesComplete exit.
}

static cl_program   MakeProgram( Type outType, Type inType, SaturationMode sat, RoundingMode round, int vectorSize, cl_kernel *outKernel )
{
    cl_program program;
    char testName[256];
    int error = 0;
    const char **strings;
    size_t stringCount = 0;

    // Create the program. This is a bit complicated because we are trying to avoid byte and short stores.
    if (0 == vectorSize)
    {
        char inName[32];
        char outName[32];
        const char *programSource[] =
        {
            "", // optional pragma
            "__kernel void ", testName, "( __global ", inName, " *src, __global ", outName, " *dest )\n"
            "{\n"
            "   size_t i = get_global_id(0);\n"
            "   dest[i] =  src[i];\n"
            "}\n"
        };
        stringCount = sizeof(programSource) / sizeof(programSource[0]);
        strings = programSource;

        if (outType == kdouble || inType == kdouble)
            programSource[0] = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";

        //create the type name
        strncpy(inName, gTypeNames[inType], sizeof(inName));
        strncpy(outName, gTypeNames[outType], sizeof(outName));
        sprintf(testName, "test_implicit_%s_%s", outName, inName);
        vlog("Building implicit %s -> %s conversion test\n", gTypeNames[inType], gTypeNames[outType]);
        fflush(stdout);
    }
    else
    {
        int vectorSizetmp = vectorSizes[vectorSize];

        char convertString[128];
        char inName[32];
        char outName[32];
        const char *programSource[] =
        {
            "", // optional pragma
            "__kernel void ", testName, "( __global ", inName, " *src, __global ", outName, " *dest )\n"
            "{\n"
            "   size_t i = get_global_id(0);\n"
            "   dest[i] = ", convertString, "( src[i] );\n"
            "}\n"
        };
        const char *programSourceV3[] =
        {
            "", // optional pragma
            "__kernel void ", testName, "( __global ", inName, " *src, __global ", outName, " *dest )\n"
            "{\n"
            "   size_t i = get_global_id(0);\n"
            "   if( i + 1 < get_global_size(0))\n"
            "       vstore3( ", convertString, "( vload3( i, src)), i, dest );\n"
            "   else\n"
            "   {\n"
            "       ", inName, "3 in;\n"
            "       ", outName, "3 out;\n"
            "       if( 0 == (i & 1) )\n"
            "           in.y = src[3*i+1];\n"
            "       in.x = src[3*i];\n"
            "       out = ", convertString, "( in ); \n"
            "       dest[3*i] = out.x;\n"
            "       if( 0 == (i & 1) )\n"
            "           dest[3*i+1] = out.y;\n"
            "   }\n"
            "}\n"
        };
        stringCount = 3 == vectorSizetmp ? sizeof(programSourceV3) / sizeof(programSourceV3[0]) :
            sizeof(programSource) / sizeof(programSource[0]);
        strings = 3 == vectorSizetmp ? programSourceV3 : programSource;

        if (outType == kdouble || inType == kdouble) {
            programSource[0] = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
            programSourceV3[0] = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
        }

        //create the type name
        switch (vectorSizetmp)
        {
        case 1:
            strncpy(inName, gTypeNames[inType], sizeof(inName));
            strncpy(outName, gTypeNames[outType], sizeof(outName));
            snprintf(convertString, sizeof(convertString), "convert_%s%s%s", outName, gSaturationNames[sat], gRoundingModeNames[round]);
            snprintf(testName, 256, "test_%s_%s", convertString, inName);
            vlog("Building %s( %s ) test\n", convertString, inName);
            break;
        case 3:
            strncpy(inName, gTypeNames[inType], sizeof(inName));
            strncpy(outName, gTypeNames[outType], sizeof(outName));
            snprintf(convertString, sizeof(convertString), "convert_%s3%s%s", outName, gSaturationNames[sat], gRoundingModeNames[round]);
            snprintf(testName, 256, "test_%s_%s3", convertString, inName);
            vlog("Building %s( %s3 ) test\n", convertString, inName);
            break;
        default:
            snprintf(inName, sizeof(inName), "%s%d", gTypeNames[inType], vectorSizetmp);
            snprintf(outName, sizeof(outName), "%s%d", gTypeNames[outType], vectorSizetmp);
            snprintf(convertString, sizeof(convertString), "convert_%s%s%s", outName, gSaturationNames[sat], gRoundingModeNames[round]);
            snprintf(testName, 256, "test_%s_%s", convertString, inName);
            vlog("Building %s( %s ) test\n", convertString, inName);
            break;
        }

        fflush(stdout);
    }
    *outKernel = NULL;

    const char *flags = NULL;
    if( gForceFTZ )
        flags = "-cl-denorms-are-zero";

    // build it
    error = create_single_kernel_helper(gContext, &program, outKernel, (cl_uint)stringCount, strings, testName, flags);
    if (error)
    {
        char    buffer[2048] = "";

        vlog_error("Failed to build kernel/program.\n", error);
        clReleaseProgram(program);
        return NULL;
    }

    return program;
}