//
// 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 <sstream>
#include <stdarg.h>
#include <stdio.h>
#include <string.h>
#if !defined(_WIN32)
#include <libgen.h>
#include <sys/mman.h>
#endif
#include <time.h>

#include <algorithm>

#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 *);

// 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, 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" );

    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;

    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__)
#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 WriteInputBufferComplete(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 / std::max(gTypeSizes[inType], gTypeSizes[outType]);
    size_t step = blockCount;
    uint64_t lastCase = 1ULL << (8 * gTypeSizes[inType]);

    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)std::min((uint64_t)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
        if ((error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_TRUE, 0,
                                          count * gTypeSizes[inType], gIn, 0,
                                          NULL, NULL)))
        {
            vlog_error("ERROR: clEnqueueWriteBuffer failed. (%d)\n", error);
            gFailCount++;
            goto exit;
        }

        // Call completion callback for the write, which will enqueue the rest
        // of the work.
        WriteInputBufferComplete((void *)&writeInputBufferInfo);

        // 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 MapResultValuesComplete(void *data);

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

    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]);

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

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

    for (vectorSize = gMinVectorSize; vectorSize < gMaxVectorSize; vectorSize++)
    {
        MapResultValuesComplete(info->calcInfo + vectorSize);
    }

    // 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 MapResultValuesComplete(void *data)
{
    cl_int status;
    CalcReferenceValuesInfo *info = (CalcReferenceValuesInfo *)data;
    cl_event calcReferenceValues = info->parent->calcReferenceValues;

    // 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;

    std::ostringstream source;
    if (outType == kdouble || inType == kdouble)
        source << "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";

    // Create the program. This is a bit complicated because we are trying to
    // avoid byte and short stores.
    if (0 == vectorSize)
    {
        // Create the type names.
        char inName[32];
        char outName[32];
        strncpy(inName, gTypeNames[inType], sizeof(inName));
        strncpy(outName, gTypeNames[outType], sizeof(outName));
        sprintf(testName, "test_implicit_%s_%s", outName, inName);

        source << "__kernel void " << testName << "( __global " << inName
               << " *src, __global " << outName << " *dest )\n";
        source << "{\n";
        source << "   size_t i = get_global_id(0);\n";
        source << "   dest[i] =  src[i];\n";
        source << "}\n";

        vlog("Building implicit %s -> %s conversion test\n", gTypeNames[inType],
             gTypeNames[outType]);
        fflush(stdout);
    }
    else
    {
        int vectorSizetmp = vectorSizes[vectorSize];

        // Create the type names.
        char convertString[128];
        char inName[32];
        char outName[32];
        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);

        if (vectorSizetmp == 3)
        {
            source << "__kernel void " << testName << "( __global " << inName
                   << " *src, __global " << outName << " *dest )\n";
            source << "{\n";
            source << "   size_t i = get_global_id(0);\n";
            source << "   if( i + 1 < get_global_size(0))\n";
            source << "       vstore3( " << convertString
                   << "( vload3( i, src)), i, dest );\n";
            source << "   else\n";
            source << "   {\n";
            source << "       " << inName << "3 in;\n";
            source << "       " << outName << "3 out;\n";
            source << "       if( 0 == (i & 1) )\n";
            source << "           in.y = src[3*i+1];\n";
            source << "       in.x = src[3*i];\n";
            source << "       out = " << convertString << "( in ); \n";
            source << "       dest[3*i] = out.x;\n";
            source << "       if( 0 == (i & 1) )\n";
            source << "           dest[3*i+1] = out.y;\n";
            source << "   }\n";
            source << "}\n";
        }
        else
        {
            source << "__kernel void " << testName << "( __global " << inName
                   << " *src, __global " << outName << " *dest )\n";
            source << "{\n";
            source << "   size_t i = get_global_id(0);\n";
            source << "   dest[i] = " << convertString << "( src[i] );\n";
            source << "}\n";
        }
    }
    *outKernel = NULL;

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

    // build it
    std::string sourceString = source.str();
    const char *programSource = sourceString.c_str();
    error = create_single_kernel_helper(gContext, &program, outKernel, 1,
                                        &programSource, testName, flags);
    if (error)
    {
        vlog_error("Failed to build kernel/program (err = %d).\n", error);
        clReleaseProgram(program);
        return NULL;
    }

    return program;
}