//
// 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 <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>

#include "procs.h"
#include "harness/conversions.h"
#include "harness/ThreadPool.h"

#define NUM_TESTS 23

#define  LONG_MATH_SHIFT_SIZE 26
#define QUICK_MATH_SHIFT_SIZE 16

static const char *kernel_code =
"__kernel void test(__global %s%s *srcA, __global %s%s *srcB, __global %s%s *dst)\n"
"{\n"
"    int  tid = get_global_id(0);\n"
"\n"
"    dst[tid] = srcA[tid] %s srcB[tid];\n"
"}\n";

static const char *kernel_code_V3 =
"__kernel void test(__global %s /*%s*/ *srcA, __global %s/*%s*/ *srcB, __global %s/*%s*/ *dst)\n"
"{\n"
"    int  tid = get_global_id(0);\n"
"\n"
"    vstore3( vload3( tid, srcA ) %s vload3( tid, srcB), tid, dst );\n"
"}\n";

static const char *kernel_code_V3_scalar_vector =
"__kernel void test(__global %s /*%s*/ *srcA, __global %s/*%s*/ *srcB, __global %s/*%s*/ *dst)\n"
"{\n"
"    int  tid = get_global_id(0);\n"
"\n"
"    vstore3( srcA[tid] %s vload3( tid, srcB), tid, dst );\n"
"}\n";

static const char *kernel_code_V3_vector_scalar =
"__kernel void test(__global %s /*%s*/ *srcA, __global %s/*%s*/ *srcB, __global %s/*%s*/ *dst)\n"
"{\n"
"    int  tid = get_global_id(0);\n"
"\n"
"    vstore3( vload3( tid, srcA ) %s srcB[tid], tid, dst );\n"
"}\n";


// Separate kernel here because it does not fit the pattern
static const char *not_kernel_code =
"__kernel void test(__global %s%s *srcA, __global %s%s *srcB, __global %s%s *dst)\n"
"{\n"
"    int  tid = get_global_id(0);\n"
"\n"
"    dst[tid] = %ssrcA[tid];\n"
"}\n";

static const char *not_kernel_code_V3 =
"__kernel void test(__global %s /*%s*/ *srcA, __global %s/*%s*/ *srcB, __global %s/*%s*/ *dst)\n"
"{\n"
"    int  tid = get_global_id(0);\n"
"\n"
"    vstore3( %s vload3( tid, srcA ), tid, dst );\n"
"}\n";

static const char *kernel_code_scalar_shift =
"__kernel void test(__global %s%s *srcA, __global %s%s *srcB, __global %s%s *dst)\n"
"{\n"
"    int  tid = get_global_id(0);\n"
"\n"
"    dst[tid] = srcA[tid] %s srcB[tid]%s;\n"
"}\n";

static const char *kernel_code_scalar_shift_V3 =
"__kernel void test(__global %s/*%s*/ *srcA, __global %s/*%s*/ *srcB, __global %s/*%s*/ *dst)\n"
"{\n"
"    int  tid = get_global_id(0);\n"
"\n"
"    vstore3( vload3( tid, srcA) %s vload3( tid, srcB )%s, tid, dst );\n"
"}\n";

static const char *kernel_code_question_colon =
"__kernel void test(__global %s%s *srcA, __global %s%s *srcB, __global %s%s *dst)\n"
"{\n"
"    int  tid = get_global_id(0);\n"
"\n"
"    dst[tid] = (srcA[tid]%s < srcB[tid]%s) ? srcA[tid] : srcB[tid];\n"
"}\n";

static const char *kernel_code_question_colon_V3 =
"__kernel void test(__global %s/*%s*/ *srcA, __global %s/*%s*/ *srcB, __global %s/*%s*/ *dst)\n"
"{\n"
"    int  tid = get_global_id(0);\n"
"\n"
"    vstore3( (vload3( tid, srcA)%s < vload3(tid, srcB)%s) ? vload3( tid, srcA) : vload3( tid, srcB), tid, dst );\n"
"}\n";


// External verification and data generation functions
extern const char *tests[];
extern const char *test_names[];
extern int verify_long(int test, size_t vector_size, cl_long *inptrA, cl_long *inptrB, cl_long *outptr, size_t n);
extern void init_long_data(uint64_t indx, int num_elements, cl_long *input_ptr[], MTdata d) ;
extern int verify_ulong(int test, size_t vector_size, cl_ulong *inptrA, cl_ulong *inptrB, cl_ulong *outptr, size_t n);
extern void init_ulong_data(uint64_t indx, int num_elements, cl_ulong *input_ptr[], MTdata d) ;
extern int verify_int(int test, size_t vector_size, cl_int *inptrA, cl_int *inptrB, cl_int *outptr, size_t n);
extern void init_int_data(uint64_t indx, int num_elements, cl_int *input_ptr[], MTdata d) ;
extern int verify_uint(int test, size_t vector_size, cl_uint *inptrA, cl_uint *inptrB, cl_uint *outptr, size_t n);
extern void init_uint_data(uint64_t indx, int num_elements, cl_uint *input_ptr[], MTdata d) ;
extern int verify_short(int test, size_t vector_size, cl_short *inptrA, cl_short *inptrB, cl_short *outptr, size_t n);
extern void init_short_data(uint64_t indx, int num_elements, cl_short *input_ptr[], MTdata d) ;
extern int verify_ushort(int test, size_t vector_size, cl_ushort *inptrA, cl_ushort *inptrB, cl_ushort *outptr, size_t n);
extern void init_ushort_data(uint64_t indx, int num_elements, cl_ushort *input_ptr[], MTdata d) ;
extern int verify_char(int test, size_t vector_size, cl_char *inptrA, cl_char *inptrB, cl_char *outptr, size_t n);
extern void init_char_data(uint64_t indx, int num_elements, cl_char *input_ptr[], MTdata d) ;
extern int verify_uchar(int test, size_t vector_size, cl_uchar *inptrA, cl_uchar *inptrB, cl_uchar *outptr, size_t n);
extern void init_uchar_data(uint64_t indx, int num_elements, cl_uchar *input_ptr[], MTdata d) ;

// Supported type list
const ExplicitType types[] = {
    kChar,
    kUChar,
    kShort,
    kUShort,
    kInt,
    kUInt,
    kLong,
    kULong,
};

enum TestStyle
{
    kDontCare=0,
    kBothVectors,
    kInputAScalar,
    kInputBScalar,
    kVectorScalarScalar,    // for the ?: operator only; indicates vector ? scalar : scalar.
    kInputCAlsoScalar = 0x80    // Or'ed flag to indicate that the selector for the ?: operator is also scalar
};

typedef struct _perThreadData
{
    cl_mem            m_streams[3];
    cl_int            *m_input_ptr[2], *m_output_ptr;
    size_t                      m_type_size;
    cl_program                m_program[NUM_TESTS];
    cl_kernel                m_kernel[NUM_TESTS];
} perThreadData;


perThreadData * perThreadDataNew()
{
    perThreadData * pThis = (perThreadData *)malloc(sizeof(perThreadData));


    memset(pThis->m_program, 0, sizeof(cl_program)*NUM_TESTS);
    memset(pThis->m_kernel, 0, sizeof(cl_kernel)*NUM_TESTS);

    pThis->m_input_ptr[0] = pThis->m_input_ptr[1] = NULL;
    pThis->m_output_ptr = NULL;

    return pThis;
}


void perThreadDataDestroy(perThreadData * pThis)
{
    int                i;
    // cleanup
    clReleaseMemObject(pThis->m_streams[0]);
    clReleaseMemObject(pThis->m_streams[1]);
    clReleaseMemObject(pThis->m_streams[2]);
    for (i=0; i<NUM_TESTS; i++)
    {
        if (pThis->m_kernel[i] != NULL) clReleaseKernel(pThis->m_kernel[i]);
        if (pThis->m_program[i] != NULL) clReleaseProgram(pThis->m_program[i]);
    }
    free(pThis->m_input_ptr[0]);
    free(pThis->m_input_ptr[1]);
    free(pThis->m_output_ptr);

    free(pThis);
}


cl_int perThreadDataInit(perThreadData * pThis, ExplicitType type,
                         int num_elements, int vectorSize,
                         int inputAVecSize, int inputBVecSize,
                         cl_context context, int start_test_ID,
                         int end_test_ID, int testID)
{
    int i;
    const char * sizeNames[] = { "", "", "2", "3", "4", "", "", "", "8", "", "", "", "", "", "", "", "16" };

    const char *type_name = get_explicit_type_name(type);
    pThis->m_type_size = get_explicit_type_size(type);
    int err;
    // Used for the && and || tests where the vector case returns a signed value
    const char *signed_type_name;
    switch (type) {
        case kChar:
        case kUChar:
            signed_type_name = get_explicit_type_name(kChar);
            break;
        case kShort:
        case kUShort:
            signed_type_name = get_explicit_type_name(kShort);
            break;
        case kInt:
        case kUInt:
            signed_type_name = get_explicit_type_name(kInt);
            break;
        case kLong:
        case kULong:
            signed_type_name = get_explicit_type_name(kLong);
            break;
        default:
            log_error("Invalid type.\n");
            return -1;
            break;
    }

    pThis->m_input_ptr[0] =
    (cl_int*)malloc(pThis->m_type_size * num_elements * vectorSize);
    pThis->m_input_ptr[1] =
    (cl_int*)malloc(pThis->m_type_size * num_elements * vectorSize);
    pThis->m_output_ptr =
    (cl_int*)malloc(pThis->m_type_size * num_elements * vectorSize);
    pThis->m_streams[0] = clCreateBuffer(
        context, CL_MEM_READ_WRITE,
        pThis->m_type_size * num_elements * inputAVecSize, NULL, &err);

    test_error(err, "clCreateBuffer failed");

    pThis->m_streams[1] = clCreateBuffer(
        context, CL_MEM_READ_WRITE,
        pThis->m_type_size * num_elements * inputBVecSize, NULL, &err);

    test_error(err, "clCreateBuffer failed");

    pThis->m_streams[2] = clCreateBuffer(
        context, CL_MEM_READ_WRITE,
        pThis->m_type_size * num_elements * vectorSize, NULL, &err);

    test_error(err, "clCreateBuffer failed");

    const char *vectorString = sizeNames[ vectorSize ];
    const char *inputAVectorString = sizeNames[ inputAVecSize ];
    const char *inputBVectorString = sizeNames[ inputBVecSize ];

    if (testID == -1)
    {
        log_info("\tTesting %s%s (%d bytes)...\n", type_name, vectorString, (int)(pThis->m_type_size*vectorSize));
    }

    char programString[4096];
    const char *ptr;


    const char * kernel_code_base = ( vectorSize != 3 ) ? kernel_code : ( inputAVecSize == 1 ) ? kernel_code_V3_scalar_vector : ( inputBVecSize == 1 ) ? kernel_code_V3_vector_scalar : kernel_code_V3;

    for (i=start_test_ID; i<end_test_ID; i++) {
        switch (i) {
            case 10:
            case 11:
                sprintf(programString, vectorSize == 3 ? kernel_code_scalar_shift_V3 : kernel_code_scalar_shift, type_name, inputAVectorString, type_name, inputBVectorString,
                        type_name, vectorString, tests[i], ((vectorSize == 1) ? "":".s0"));
                break;
            case 12:
                sprintf(programString, vectorSize == 3 ? not_kernel_code_V3 : not_kernel_code, type_name, inputAVectorString, type_name, inputBVectorString,
                        type_name, vectorString, tests[i]);
                break;
            case 13:
                sprintf(programString, vectorSize == 3 ? kernel_code_question_colon_V3 : kernel_code_question_colon,
                        type_name, inputAVectorString, type_name, inputBVectorString,
                        type_name, vectorString, ((vectorSize == 1) ? "":".s0"), ((vectorSize == 1) ? "":".s0")) ;
                break;
            case 14:
            case 15:
            case 16:
            case 17:
            case 18:
            case 19:
            case 20:
            case 21:
                // Need an unsigned result here for vector sizes > 1
                sprintf(programString, kernel_code_base, type_name, inputAVectorString, type_name, inputBVectorString,
                        ((vectorSize == 1) ? type_name : signed_type_name), vectorString, tests[i]);
                break;
            case 22:
                // Need an unsigned result here for vector sizes > 1
                sprintf(programString, vectorSize == 3 ? not_kernel_code_V3 : not_kernel_code, type_name, inputAVectorString, type_name, inputBVectorString,
                        ((vectorSize == 1) ? type_name : signed_type_name), vectorString, tests[i]);
                break;
            default:
                sprintf(programString, kernel_code_base, type_name, inputAVectorString, type_name, inputBVectorString,
                        type_name, vectorString, tests[i]);
                break;
        }

        //printf("kernel: %s\n", programString);
        ptr = programString;
        err = create_single_kernel_helper( context,
                                          &(pThis->m_program[ i ]),
                                          &(pThis->m_kernel[ i ]), 1,
                                          &ptr, "test" );
        test_error( err, "Unable to create test kernel" );
        err = clSetKernelArg(pThis->m_kernel[i], 0,
                             sizeof pThis->m_streams[0],
                             &(pThis->m_streams[0]) );
        err |= clSetKernelArg(pThis->m_kernel[i], 1,
                              sizeof pThis->m_streams[1],
                              &(pThis->m_streams[1]) );
        err |= clSetKernelArg(pThis->m_kernel[i], 2,
                              sizeof pThis->m_streams[2],
                              &(pThis->m_streams[2]) );
        test_error(err, "clSetKernelArgs failed");
    }

    return CL_SUCCESS;
}

typedef struct _globalThreadData
{
    cl_device_id     m_deviceID;
    cl_context       m_context;
    // cl_command_queue m_queue;
    int              m_num_elements;
    int              m_threadcount;
    int              m_vectorSize;
    int              m_num_runs_shift;
    TestStyle        m_style;
    ExplicitType     m_type;
    MTdata *         m_pRandData;
    uint64_t         m_offset;
    int              m_testID;
    perThreadData  **m_arrPerThreadData;
} globalThreadData;


globalThreadData * globalThreadDataNew(cl_device_id deviceID, cl_context context,
                                       cl_command_queue queue, int num_elements,
                                       int vectorSize, TestStyle style, int num_runs_shift,
                                       ExplicitType type, int testID,
                                       int threadcount)
{
    int i;
    globalThreadData * pThis = (globalThreadData *)malloc(sizeof(globalThreadData));
    pThis->m_deviceID = deviceID;
    pThis->m_context = context;
    // pThis->m_queue = queue;
    pThis->m_num_elements = num_elements;
    pThis->m_num_runs_shift = num_runs_shift;
    pThis->m_vectorSize = vectorSize;
    pThis->m_style = style;
    pThis->m_type = type;
    pThis->m_offset = (uint64_t)0;
    pThis->m_testID = testID;
    pThis->m_arrPerThreadData = NULL;
    pThis->m_threadcount = threadcount;

    pThis->m_pRandData = (MTdata *)malloc(threadcount*sizeof(MTdata));
    pThis->m_arrPerThreadData = (perThreadData **)
    malloc(threadcount*sizeof(perThreadData *));
    for(i=0; i < threadcount; ++i)
    {
        pThis->m_pRandData[i] = init_genrand(i+1);
        pThis->m_arrPerThreadData[i] = NULL;
    }

    return pThis;
}

void globalThreadDataDestroy(globalThreadData * pThis)
{
    int i;

    for(i=0; i < pThis->m_threadcount; ++i)
    {
        free_mtdata(pThis->m_pRandData[i]);
        if(pThis->m_arrPerThreadData[i] != NULL)
        {
            perThreadDataDestroy(pThis->m_arrPerThreadData[i]);
        }
    }
    free(pThis->m_arrPerThreadData);
    free(pThis->m_pRandData);
    free(pThis);
}

int
test_integer_ops(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, int vectorSize, TestStyle style, int num_runs_shift, ExplicitType type, int testID, MTdata randIn, uint64_t startIndx, uint64_t endIndx,
                 perThreadData ** ppThreadData);


cl_int test_integer_ops_do_thread( cl_uint job_id, cl_uint thread_id, void *userInfo )
{
    cl_int error; cl_int result;
    globalThreadData * threadInfoGlobal = (globalThreadData *)userInfo;
    cl_command_queue queue;

#if THREAD_DEBUG
    log_error("Thread %x (job %x) about to create command queue\n",
              thread_id, job_id);
#endif

    queue =  clCreateCommandQueue (threadInfoGlobal->m_context,
                                   threadInfoGlobal->m_deviceID,0,
                                   &error);

    if(error != CL_SUCCESS)
    {
        log_error("Thread %x (job %x) could not create command queue\n",
                  thread_id, job_id);
        return error; // should we clean up the queue too?
    }

#if THREAD_DEBUG
    log_error("Thread %x (job %x) created command queue\n",
              thread_id, job_id);
#endif

    result = test_integer_ops(  threadInfoGlobal->m_deviceID,
                              threadInfoGlobal->m_context,
                              queue,
                              threadInfoGlobal->m_num_elements,
                              threadInfoGlobal->m_vectorSize, threadInfoGlobal->m_style,
                              threadInfoGlobal->m_num_runs_shift,
                              threadInfoGlobal->m_type, threadInfoGlobal->m_testID,
                              threadInfoGlobal->m_pRandData[thread_id],
                              threadInfoGlobal->m_offset + threadInfoGlobal->m_num_elements*job_id,
                              threadInfoGlobal->m_offset + threadInfoGlobal->m_num_elements*(job_id+1),
                              &(threadInfoGlobal->m_arrPerThreadData[thread_id])
                              );

    if(result != 0)
    {
        log_error("Thread %x (job %x) failed test_integer_ops with result %x\n",
                  thread_id, job_id, result);
        // return error;
    }


    error = clReleaseCommandQueue(queue);
    if(error != CL_SUCCESS)
    {
        log_error("Thread %x (job %x) could not release command queue\n",
                  thread_id, job_id);
        return error;
    }
    return result;
}

int
test_integer_ops_threaded(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, int vectorSize, TestStyle style, int num_runs_shift, ExplicitType type, int testID)
{
    globalThreadData * pThreadInfo = NULL;
    cl_int result=0;
    cl_uint threadcount = GetThreadCount();

  // Check to see if we are using single threaded mode on other than a 1.0 device
  if (getenv( "CL_TEST_SINGLE_THREADED" )) {

    char device_version[1024] = { 0 };
    result = clGetDeviceInfo( deviceID, CL_DEVICE_VERSION, sizeof(device_version), device_version, NULL );
    if(result != CL_SUCCESS)
    {
      log_error("clGetDeviceInfo(CL_DEVICE_GLOBAL_MEM_SIZE) failed: %d\n", result);
      return result;
    }

    if (strcmp("OpenCL 1.0 ",device_version)) {
      log_error("ERROR: CL_TEST_SINGLE_THREADED is set in the environment. Running single threaded.\n");
    }
  }

    // This test will run threadcount threads concurrently; each thread will execute test_integer_ops()
    // which will allocate 2 OpenCL buffers on the device; each buffer has size num_elements * type_size * vectorSize.
    // We need to make sure that the total device memory allocated by all threads does not exceed the maximum
    // memory on the device. If it does, we decrease num_elements until all threads combined will not
    // over-subscribe device memory.
    cl_ulong maxDeviceGlobalMem;
    result = clGetDeviceInfo(deviceID, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(maxDeviceGlobalMem), &maxDeviceGlobalMem, NULL);
    if(result != CL_SUCCESS)
    {
        log_error("clGetDeviceInfo(CL_DEVICE_GLOBAL_MEM_SIZE) failed: %d\n", result);
        return result;
    }

  if (maxDeviceGlobalMem > (cl_ulong)SIZE_MAX) {
    maxDeviceGlobalMem = (cl_ulong)SIZE_MAX;
  }

    // Let's not take all device memory - reduce by 75%
    maxDeviceGlobalMem = (maxDeviceGlobalMem * 3) >> 2;
    // Now reduce num_elements so that the total device memory usage does not exceed 75% of global device memory.
    size_t type_size = get_explicit_type_size(type);
    while ((cl_ulong)threadcount * 4 * num_elements * type_size * vectorSize > maxDeviceGlobalMem)
    {
        num_elements >>= 1;
    }

    uint64_t startIndx = (uint64_t)0;
    uint64_t endIndx = (1ULL<<num_runs_shift);
    uint64_t jobcount = (endIndx-startIndx)/num_elements;

    if(jobcount==0)
    {
        jobcount = 1;
    }

    pThreadInfo = globalThreadDataNew(deviceID, context, queue, num_elements,
                                      vectorSize, style, num_runs_shift,
                                      type, testID, threadcount);


    pThreadInfo->m_offset = startIndx;

#if THREAD_DEBUG
    log_error("Launching %llx jobs\n",
              jobcount);
#endif

    result = ThreadPool_Do(test_integer_ops_do_thread, (cl_uint)jobcount, (void *)pThreadInfo);

    if(result != 0)
    {
        // cleanup ??
        log_error("ThreadPool_Do return non-success value %d\n", result);

    }
    globalThreadDataDestroy(pThreadInfo);
    return result;
}


int
test_integer_ops(cl_device_id deviceID, cl_context context,
                 cl_command_queue queue, int num_elements,
                 int vectorSize, TestStyle style, int num_runs_shift,
                 ExplicitType type, int testID, MTdata randDataIn,
                 uint64_t startIndx, uint64_t endIndx,
                 perThreadData ** ppThreadData)
{
    size_t    threads[1];
    int                err;
    int                i;
    int inputAVecSize, inputBVecSize;


    inputAVecSize = inputBVecSize = vectorSize;
    if( style == kInputAScalar )
        inputAVecSize = 1;
    else if( style == kInputBScalar )
        inputBVecSize = 1;

    /*
     if( inputAVecSize != inputBVecSize )
     log_info("Testing \"%s\" on %s%d (%s-%s inputs) (range %llx - %llx of 0-%llx)\n",
     test_names[testID],
     get_explicit_type_name(type), vectorSize,
     ( inputAVecSize == 1 ) ? "scalar" : "vector",
     ( inputBVecSize == 1 ) ? "scalar" : "vector",
     startIndx, endIndx, (1ULL<<num_runs_shift) );
     else
     log_info("Testing \"%s\" on %s%d (range %llx - %llx of 0-%llx)\n",
     test_names[testID],
     get_explicit_type_name(type), vectorSize,
     startIndx, endIndx, (1ULL<<num_runs_shift));
     */


    // Figure out which sub-test to run, or all of them
    int start_test_ID = 0;
    int end_test_ID = NUM_TESTS;
    if (testID != -1) {
        start_test_ID = testID;
        end_test_ID = testID+1;
    }
    if (testID > NUM_TESTS) {
        log_error("Invalid test ID: %d\n", testID);
        return -1;
    }

    if(*ppThreadData == NULL)
    {
        *ppThreadData = perThreadDataNew();
        err = perThreadDataInit(*ppThreadData,
                                type, num_elements, vectorSize,
                                inputAVecSize, inputBVecSize,
                                context, start_test_ID,
                                end_test_ID, testID);
        test_error(err, "failed to init per thread data\n");
    }

    perThreadData * pThreadData = *ppThreadData;


    threads[0] = (size_t)num_elements;
    int error_count = 0;
    for (i=start_test_ID; i<end_test_ID; i++)
    {
        uint64_t    indx;


        if(startIndx >= endIndx)
        {
            startIndx = (uint64_t)0;
            endIndx = (1ULL<<num_runs_shift);
        }
        for (indx=startIndx; indx < endIndx; indx+=num_elements)
        {

            switch (type) {
                case     kChar:
                    init_char_data(indx, num_elements * vectorSize, (cl_char**)(pThreadData->m_input_ptr), randDataIn);
                    break;
                case     kUChar:
                    init_uchar_data(indx, num_elements * vectorSize, (cl_uchar**)(pThreadData->m_input_ptr), randDataIn);
                    break;
                case     kShort:
                    init_short_data(indx, num_elements * vectorSize, (cl_short**)(pThreadData->m_input_ptr), randDataIn);
                    break;
                case     kUShort:
                    init_ushort_data(indx, num_elements * vectorSize, (cl_ushort**)(pThreadData->m_input_ptr), randDataIn);
                    break;
                case     kInt:
                    init_int_data(indx, num_elements * vectorSize, (cl_int**)(pThreadData->m_input_ptr), randDataIn);
                    break;
                case     kUInt:
                    init_uint_data(indx, num_elements * vectorSize, (cl_uint**)(pThreadData->m_input_ptr), randDataIn);
                    break;
                case     kLong:
                    init_long_data(indx, num_elements * vectorSize, (cl_long**)(pThreadData->m_input_ptr), randDataIn);
                    break;
                case     kULong:
                    init_ulong_data(indx, num_elements * vectorSize, (cl_ulong**)(pThreadData->m_input_ptr), randDataIn);
                    break;
                default:
                    err = 1;
                    log_error("Invalid type.\n");
                    break;
            }


            err = clEnqueueWriteBuffer(queue, pThreadData->m_streams[0], CL_FALSE, 0, pThreadData->m_type_size*num_elements * inputAVecSize, (void *)pThreadData->m_input_ptr[0], 0, NULL, NULL);
            test_error(err, "clEnqueueWriteBuffer failed");
            err = clEnqueueWriteBuffer( queue, pThreadData->m_streams[1], CL_FALSE, 0, pThreadData->m_type_size*num_elements * inputBVecSize, (void *)pThreadData->m_input_ptr[1], 0, NULL, NULL );
            test_error(err, "clEnqueueWriteBuffer failed");

            err = clEnqueueNDRangeKernel( queue, pThreadData->m_kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL );
            test_error(err, "clEnqueueNDRangeKernel failed");

            err = clEnqueueReadBuffer( queue, pThreadData->m_streams[2], CL_TRUE, 0, pThreadData->m_type_size*num_elements * vectorSize, (void *)pThreadData->m_output_ptr, 0, NULL, NULL );
            test_error(err, "clEnqueueReadBuffer failed");

            // log_info("Performing verification\n");

            // If one of the inputs are scalar, we need to extend the input values to vectors
            // to accommodate the verify functions
            if( vectorSize > 1 )
            {
                char * p = NULL;
                if( style == kInputAScalar )
                    p = (char *)pThreadData->m_input_ptr[ 0 ];
                else if( style == kInputBScalar )
                    p = (char *)pThreadData->m_input_ptr[ 1 ];
                if( p != NULL )
                {
                    for( int element = num_elements - 1; element >= 0; element-- )
                    {
                        for( int vec = ( element == 0 ) ? 1 : 0; vec < vectorSize; vec++ )
                            memcpy( p + ( element * vectorSize + vec ) * pThreadData->m_type_size, p + element * pThreadData->m_type_size, pThreadData->m_type_size );
                    }
                }
            }

            switch (type) {
                case     kChar:
                    err = verify_char(i, vectorSize, (cl_char*)pThreadData->m_input_ptr[0], (cl_char*)pThreadData->m_input_ptr[1], (cl_char*)pThreadData->m_output_ptr, num_elements * vectorSize);
                    break;
                case     kUChar:
                    err = verify_uchar(i, vectorSize, (cl_uchar*)pThreadData->m_input_ptr[0], (cl_uchar*)pThreadData->m_input_ptr[1], (cl_uchar*)pThreadData->m_output_ptr, num_elements * vectorSize);
                    break;
                case     kShort:
                    err = verify_short(i, vectorSize, (cl_short*)pThreadData->m_input_ptr[0], (cl_short*)pThreadData->m_input_ptr[1], (cl_short*)pThreadData->m_output_ptr, num_elements * vectorSize);
                    break;
                case     kUShort:
                    err = verify_ushort(i, vectorSize, (cl_ushort*)pThreadData->m_input_ptr[0], (cl_ushort*)pThreadData->m_input_ptr[1], (cl_ushort*)pThreadData->m_output_ptr, num_elements * vectorSize);
                    break;
                case     kInt:
                    err = verify_int(i, vectorSize, (cl_int*)pThreadData->m_input_ptr[0], (cl_int*)pThreadData->m_input_ptr[1], (cl_int*)pThreadData->m_output_ptr, num_elements * vectorSize);
                    break;
                case     kUInt:
                    err = verify_uint(i, vectorSize, (cl_uint*)pThreadData->m_input_ptr[0], (cl_uint*)pThreadData->m_input_ptr[1], (cl_uint*)pThreadData->m_output_ptr, num_elements * vectorSize);
                    break;
                case     kLong:
                    err = verify_long(i, vectorSize, (cl_long*)pThreadData->m_input_ptr[0], (cl_long*)pThreadData->m_input_ptr[1], (cl_long*)pThreadData->m_output_ptr, num_elements * vectorSize);
                    break;
                case     kULong:
                    err = verify_ulong(i, vectorSize, (cl_ulong*)pThreadData->m_input_ptr[0], (cl_ulong*)pThreadData->m_input_ptr[1], (cl_ulong*)pThreadData->m_output_ptr, num_elements * vectorSize);
                    break;
                default:
                    err = 1;
                    log_error("Invalid type.\n");
                    break;
            }

            if (err) {
#if 0
                log_error( "* inASize: %d inBSize: %d numElem: %d\n", inputAVecSize, inputBVecSize, num_elements );
                cl_char *inP = (cl_char *)pThreadData->m_input_ptr[0];
                log_error( "from 18:\n" );
                for( int q = 18; q < 64; q++ )
                {
                    log_error( "%02x ", inP[ q ] );
                }
                log_error( "\n" );
                inP = (cl_char *)pThreadData->m_input_ptr[1];
                for( int q = 18; q < 64; q++ )
                {
                    log_error( "%02x ", inP[ q ] );
                }
                log_error( "\n" );
                inP = (cl_char *)pThreadData->m_output_ptr;
                for( int q = 18; q < 64; q++ )
                {
                    log_error( "%02x ", inP[ q ] );
                }
                log_error( "\n" );
                log_error( "from 36:\n" );
                inP = (cl_char *)pThreadData->m_input_ptr[0];
                for( int q = 36; q < 64; q++ )
                {
                    log_error( "%02x ", inP[ q ] );
                }
                log_error( "\n" );
                inP = (cl_char *)pThreadData->m_input_ptr[1];
                for( int q = 36; q < 64; q++ )
                {
                    log_error( "%02x ", inP[ q ] );
                }
                log_error( "\n" );
                inP = (cl_char *)pThreadData->m_output_ptr;
                for( int q = 36; q < 64; q++ )
                {
                    log_error( "%02x ", inP[ q ] );
                }
                log_error( "\n" );
#endif
                error_count++;
                break;
            }
        }

        /*

         const char * sizeNames[] = { "", "", "2", "3", "4", "", "", "", "8", "", "", "", "", "", "", "", "16" };

         if (err) {
         log_error("\t\t%s%s test %s failed (range %llx - %llx of 0-%llx)\n",
         get_explicit_type_name(type), sizeNames[vectorSize],
         test_names[i],
         startIndx, endIndx,
         (1ULL<<num_runs_shift));
         } else {
         log_info("\t\t%s%s test %s passed (range %llx - %llx of 0-%llx)\n",
         get_explicit_type_name(type), sizeNames[vectorSize],
         test_names[i],
         startIndx, endIndx,
         (1ULL<<num_runs_shift));
         }
         */
    }


    return error_count;
}


// Run all the vector sizes for a given test
int run_specific_test(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, ExplicitType type, int num, int testID) {
    int errors = 0;
    errors += test_integer_ops_threaded(deviceID, context, queue, (1024*1024*2)/1, 1, kBothVectors, num, type, testID);
    errors += test_integer_ops_threaded(deviceID, context, queue, (1024*1024*2)/2, 2, kBothVectors, num, type, testID);
    errors += test_integer_ops_threaded(deviceID, context, queue, (1024*1024*2)/3, 3, kBothVectors, num, type, testID);
    errors += test_integer_ops_threaded(deviceID, context, queue, (1024*1024*2)/4, 4, kBothVectors, num, type, testID);
    errors += test_integer_ops_threaded(deviceID, context, queue, (1024*1024*2)/8, 8, kBothVectors, num, type, testID);
    errors += test_integer_ops_threaded(deviceID, context, queue, (1024*1024*2)/16, 16, kBothVectors, num, type, testID);
    return errors;
}

// Run multiple tests for a given type
int run_multiple_tests(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, ExplicitType type, int num, int *tests, int total_tests) {
    int errors = 0;

    if (getenv("CL_WIMPY_MODE") && num == LONG_MATH_SHIFT_SIZE) {
      log_info("Detected CL_WIMPY_MODE env\n");
      log_info("Skipping long test\n");
      return 0;
    }

    int i;
    for (i=0; i<total_tests; i++)
    {
        int localErrors;
        log_info("Testing \"%s\" ", test_names[tests[i]]);  fflush( stdout );
        localErrors = run_specific_test(deviceID, context, queue, num_elements, type, num, tests[i]);
        if( localErrors )
            log_info( "FAILED\n" );
        else
            log_info( "passed\n" );

        errors += localErrors;
    }

    return errors;
}

// Run all the math tests for a given type
int run_test_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, ExplicitType type, int num) {
    int tests[] = {0, 1, 2, 3, 4};
    return run_multiple_tests(deviceID, context, queue, num_elements, type, num, tests, (int)(sizeof(tests)/sizeof(int)));
}

// Run all the logic tests for a given type
int run_test_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, ExplicitType type, int num) {
    int tests[] = {5, 6, 7, 12, 14, 15, 22};
    return run_multiple_tests(deviceID, context, queue, num_elements, type, num, tests, (int)(sizeof(tests)/sizeof(int)));
}

// Run all the shifting tests for a given type
int run_test_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, ExplicitType type, int num) {
    int tests[] = {8, 9, 10, 11};
    return run_multiple_tests(deviceID, context, queue, num_elements, type, num, tests, (int)(sizeof(tests)/sizeof(int)));
}

// Run all the comparison tests for a given type
int run_test_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, ExplicitType type, int num) {
    int tests[] = {13, 16, 17, 18, 19, 20, 21};
    return run_multiple_tests(deviceID, context, queue, num_elements, type, num, tests, (int)(sizeof(tests)/sizeof(int)));
}

// Run all tests for a given type
int run_test(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, ExplicitType type, int num) {
    int errors = 0;
    errors += test_integer_ops_threaded(deviceID, context, queue, 1024*1024*2, 1, kBothVectors, num, type, -1);
    errors += test_integer_ops_threaded(deviceID, context, queue, 1024*1024*2, 2, kBothVectors, num, type, -1);
    errors += test_integer_ops_threaded(deviceID, context, queue, 1024*1024*2, 3, kBothVectors, num, type, -1);
    errors += test_integer_ops_threaded(deviceID, context, queue, 1024*1024*2, 4, kBothVectors, num, type, -1);
    errors += test_integer_ops_threaded(deviceID, context, queue, 1024*1024*2, 8, kBothVectors, num, type, -1);
    errors += test_integer_ops_threaded(deviceID, context, queue, 1024*1024*2, 16, kBothVectors, num, type, -1);
    return errors;
}


// -----------------
// Long tests
// -----------------
int test_long_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_math(deviceID, context, queue, num_elements, kLong, LONG_MATH_SHIFT_SIZE);
}
int test_long_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_logic(deviceID, context, queue, num_elements, kLong, LONG_MATH_SHIFT_SIZE);
}
int test_long_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_shift(deviceID, context, queue, num_elements, kLong, LONG_MATH_SHIFT_SIZE);
}
int test_long_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_compare(deviceID, context, queue, num_elements, kLong, LONG_MATH_SHIFT_SIZE);
}
int test_quick_long_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_math(deviceID, context, queue, num_elements, kLong, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_long_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_logic(deviceID, context, queue, num_elements, kLong, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_long_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_shift(deviceID, context, queue, num_elements, kLong, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_long_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_compare(deviceID, context, queue, num_elements, kLong, QUICK_MATH_SHIFT_SIZE);
}


// -----------------
// ULong tests
// -----------------
int test_ulong_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_math(deviceID, context, queue, num_elements, kULong, LONG_MATH_SHIFT_SIZE);
}
int test_ulong_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_logic(deviceID, context, queue, num_elements, kULong, LONG_MATH_SHIFT_SIZE);
}
int test_ulong_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_shift(deviceID, context, queue, num_elements, kULong, LONG_MATH_SHIFT_SIZE);
}
int test_ulong_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_compare(deviceID, context, queue, num_elements, kULong, LONG_MATH_SHIFT_SIZE);
}
int test_quick_ulong_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_math(deviceID, context, queue, num_elements, kULong, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_ulong_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_logic(deviceID, context, queue, num_elements, kULong, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_ulong_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_shift(deviceID, context, queue, num_elements, kULong, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_ulong_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test_compare(deviceID, context, queue, num_elements, kULong, QUICK_MATH_SHIFT_SIZE);
}


// -----------------
// Int tests
// -----------------
int test_int_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_math(deviceID, context, queue, num_elements, kInt, LONG_MATH_SHIFT_SIZE);
}
int test_int_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_logic(deviceID, context, queue, num_elements, kInt, LONG_MATH_SHIFT_SIZE);
}
int test_int_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_shift(deviceID, context, queue, num_elements, kInt, LONG_MATH_SHIFT_SIZE);
}
int test_int_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_compare(deviceID, context, queue, num_elements, kInt, LONG_MATH_SHIFT_SIZE);
}
int test_quick_int_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_math(deviceID, context, queue, num_elements, kInt, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_int_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_logic(deviceID, context, queue, num_elements, kInt, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_int_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_shift(deviceID, context, queue, num_elements, kInt, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_int_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_compare(deviceID, context, queue, num_elements, kInt, QUICK_MATH_SHIFT_SIZE);
}


// -----------------
// UInt tests
// -----------------
int test_uint_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_math(deviceID, context, queue, num_elements, kUInt, LONG_MATH_SHIFT_SIZE);
}
int test_uint_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_logic(deviceID, context, queue, num_elements, kUInt, LONG_MATH_SHIFT_SIZE);
}
int test_uint_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_shift(deviceID, context, queue, num_elements, kUInt, LONG_MATH_SHIFT_SIZE);
}
int test_uint_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_compare(deviceID, context, queue, num_elements, kUInt, LONG_MATH_SHIFT_SIZE);
}
int test_quick_uint_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_math(deviceID, context, queue, num_elements, kUInt, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_uint_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_logic(deviceID, context, queue, num_elements, kUInt, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_uint_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_shift(deviceID, context, queue, num_elements, kUInt, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_uint_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_compare(deviceID, context, queue, num_elements, kUInt, QUICK_MATH_SHIFT_SIZE);
}


// -----------------
// Short tests
// -----------------
int test_short_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_math(deviceID, context, queue, num_elements, kShort, LONG_MATH_SHIFT_SIZE);
}
int test_short_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_logic(deviceID, context, queue, num_elements, kShort, LONG_MATH_SHIFT_SIZE);
}
int test_short_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_shift(deviceID, context, queue, num_elements, kShort, LONG_MATH_SHIFT_SIZE);
}
int test_short_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_compare(deviceID, context, queue, num_elements, kShort, LONG_MATH_SHIFT_SIZE);
}
int test_quick_short_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_math(deviceID, context, queue, num_elements, kShort, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_short_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_logic(deviceID, context, queue, num_elements, kShort, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_short_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_shift(deviceID, context, queue, num_elements, kShort, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_short_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_compare(deviceID, context, queue, num_elements, kShort, QUICK_MATH_SHIFT_SIZE);
}


// -----------------
// UShort tests
// -----------------
int test_ushort_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_math(deviceID, context, queue, num_elements, kUShort, LONG_MATH_SHIFT_SIZE);
}
int test_ushort_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_logic(deviceID, context, queue, num_elements, kUShort, LONG_MATH_SHIFT_SIZE);
}
int test_ushort_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_shift(deviceID, context, queue, num_elements, kUShort, LONG_MATH_SHIFT_SIZE);
}
int test_ushort_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_compare(deviceID, context, queue, num_elements, kUShort, LONG_MATH_SHIFT_SIZE);
}
int test_quick_ushort_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_math(deviceID, context, queue, num_elements, kUShort, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_ushort_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_logic(deviceID, context, queue, num_elements, kUShort, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_ushort_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_shift(deviceID, context, queue, num_elements, kUShort, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_ushort_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_compare(deviceID, context, queue, num_elements, kUShort, QUICK_MATH_SHIFT_SIZE);
}


// -----------------
// Char tests
// -----------------
int test_char_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_math(deviceID, context, queue, num_elements, kChar, LONG_MATH_SHIFT_SIZE);
}
int test_char_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_logic(deviceID, context, queue, num_elements, kChar, LONG_MATH_SHIFT_SIZE);
}
int test_char_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_shift(deviceID, context, queue, num_elements, kChar, LONG_MATH_SHIFT_SIZE);
}
int test_char_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_compare(deviceID, context, queue, num_elements, kChar, LONG_MATH_SHIFT_SIZE);
}
int test_quick_char_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_math(deviceID, context, queue, num_elements, kChar, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_char_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_logic(deviceID, context, queue, num_elements, kChar, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_char_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_shift(deviceID, context, queue, num_elements, kChar, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_char_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_compare(deviceID, context, queue, num_elements, kChar, QUICK_MATH_SHIFT_SIZE);
}


// -----------------
// UChar tests
// -----------------
int test_uchar_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_math(deviceID, context, queue, num_elements, kUChar, LONG_MATH_SHIFT_SIZE);
}
int test_uchar_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_logic(deviceID, context, queue, num_elements, kUChar, LONG_MATH_SHIFT_SIZE);
}
int test_uchar_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_shift(deviceID, context, queue, num_elements, kUChar, LONG_MATH_SHIFT_SIZE);
}
int test_uchar_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_compare(deviceID, context, queue, num_elements, kUChar, LONG_MATH_SHIFT_SIZE);
}
int test_quick_uchar_math(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_math(deviceID, context, queue, num_elements, kUChar, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_uchar_logic(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_logic(deviceID, context, queue, num_elements, kUChar, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_uchar_shift(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_shift(deviceID, context, queue, num_elements, kUChar, QUICK_MATH_SHIFT_SIZE);
}
int test_quick_uchar_compare(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test_compare(deviceID, context, queue, num_elements, kUChar, QUICK_MATH_SHIFT_SIZE);
}


// These are kept for debugging if you want to run all the tests together.

int test_long(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test(deviceID, context, queue, num_elements, kLong, LONG_MATH_SHIFT_SIZE);
}

int test_quick_long(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test(deviceID, context, queue, num_elements, kLong, QUICK_MATH_SHIFT_SIZE);
}

int test_ulong(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test(deviceID, context, queue, num_elements, kULong, LONG_MATH_SHIFT_SIZE);
}

int test_quick_ulong(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    if (!gHasLong)
    {
        log_info( "WARNING: 64 bit integers are not supported on this device. Skipping\n" );
        return CL_SUCCESS;
    }
    return run_test(deviceID, context, queue, num_elements, kULong, QUICK_MATH_SHIFT_SIZE);
}

int test_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test(deviceID, context, queue, num_elements, kInt, LONG_MATH_SHIFT_SIZE);
}

int test_quick_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test(deviceID, context, queue, num_elements, kInt, QUICK_MATH_SHIFT_SIZE);
}

int test_uint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test(deviceID, context, queue, num_elements, kUInt, LONG_MATH_SHIFT_SIZE);
}

int test_quick_uint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test(deviceID, context, queue, num_elements, kUInt, QUICK_MATH_SHIFT_SIZE);
}

int test_short(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test(deviceID, context, queue, num_elements, kShort, LONG_MATH_SHIFT_SIZE);
}

int test_quick_short(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test(deviceID, context, queue, num_elements, kShort, QUICK_MATH_SHIFT_SIZE);
}

int test_ushort(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test(deviceID, context, queue, num_elements, kUShort, LONG_MATH_SHIFT_SIZE);
}

int test_quick_ushort(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test(deviceID, context, queue, num_elements, kUShort, QUICK_MATH_SHIFT_SIZE);
}

int test_char(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test(deviceID, context, queue, num_elements, kChar, LONG_MATH_SHIFT_SIZE);
}

int test_quick_char(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test(deviceID, context, queue, num_elements, kChar, QUICK_MATH_SHIFT_SIZE);
}

int test_uchar(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test(deviceID, context, queue, num_elements, kUChar, LONG_MATH_SHIFT_SIZE);
}

int test_quick_uchar(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) {
    return run_test(deviceID, context, queue, num_elements, kUChar, QUICK_MATH_SHIFT_SIZE);
}

// Prototype for below
int test_question_colon_op(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements,
                           int vectorSize, TestStyle style, ExplicitType type );

// Run all the vector sizes for a given test in scalar-vector and vector-scalar modes
int run_test_sizes(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, ExplicitType type, int num, int testID)
{
    int sizes[] = { 2, 3, 4, 8, 16, 0 };
    int errors = 0;

    for( int i = 0; sizes[ i ] != 0; i++ )
    {
        if( testID == 13 )
        {
            errors += test_question_colon_op( deviceID, context, queue, num_elements / sizes[i], sizes[i], kInputAScalar, type );
            errors += test_question_colon_op( deviceID, context, queue, num_elements / sizes[i], sizes[i], kInputBScalar, type );
            errors += test_question_colon_op( deviceID, context, queue, num_elements / sizes[i], sizes[i], kVectorScalarScalar, type );

            errors += test_question_colon_op( deviceID, context, queue, num_elements / sizes[i], sizes[i], (TestStyle)(kBothVectors | kInputCAlsoScalar), type );
            errors += test_question_colon_op( deviceID, context, queue, num_elements / sizes[i], sizes[i], (TestStyle)(kInputAScalar | kInputCAlsoScalar), type );
            errors += test_question_colon_op( deviceID, context, queue, num_elements / sizes[i], sizes[i], (TestStyle)(kInputBScalar | kInputCAlsoScalar), type );
        }
        else
        {
            errors += test_integer_ops_threaded(deviceID, context, queue, num_elements / sizes[i], sizes[i], kInputAScalar, num, type, testID);
            errors += test_integer_ops_threaded(deviceID, context, queue, num_elements / sizes[i], sizes[i], kInputBScalar, num, type, testID);
        }
    }
    return errors;
}

// Run all the tests for scalar-vector and vector-scalar for a given type
int run_vector_scalar_tests( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, ExplicitType type, int num )
{
    int errors = 0;
    size_t i;

    // Shift operators:
    // a) cannot take scalars as first parameter and vectors as second
    // b) have the vector >> scalar case tested by tests 10 and 11
    // so they get skipped entirely

    int testsToRun[] = { 0, 1, 2, 3, 4, 5, 6, 7,
        13, 14, 15, 16, 17, 18, 19, 20, 21 };
    for (i=0; i< sizeof(testsToRun)/sizeof(testsToRun[0]); i++)
    {
        errors += run_test_sizes(deviceID, context, queue, 2048, type, num, testsToRun[i]);
    }
    return errors;
}

int test_vector_scalar(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
    int errors = 0;
    int numTypes = sizeof( types ) / sizeof( types[ 0 ] );

    for( int t = 0; t < numTypes; t++ )
    {
        if ((types[ t ] == kLong || types[ t ] == kULong) && !gHasLong)
            continue;

        errors += run_vector_scalar_tests( deviceID, context, queue, num_elements, types[ t ], 1 );
        break;
    }

    return errors;
}

void generate_random_bool_data( size_t count, MTdata d, cl_char *outData, size_t outDataSize )
{
    cl_uint bits = genrand_int32(d);
    cl_uint bitsLeft = 32;

    memset( outData, 0, outDataSize * count );

    for( size_t i = 0; i < count; i++ )
    {
        if( 0 == bitsLeft)
        {
            bits = genrand_int32(d);
            bitsLeft = 32;
        }

        // Note: we will be setting just any bit non-zero for the type, so we can easily skip past
        // and just write bytes (assuming the entire output buffer is already zeroed, which we did)
        *outData = ( bits & 1 ) ? 0xff : 0;

        bits >>= 1; bitsLeft -= 1;

        outData += outDataSize;
    }
}

static const char *kernel_question_colon_full =
"__kernel void test(__global %s%s *srcA, __global %s%s *srcB, __global %s%s *srcC, __global %s%s *dst)\n"
"{\n"
"    int  tid = get_global_id(0);\n"
"\n"
"    %s%s valA = %ssrcA%s"
"    %s%s valB = %ssrcB%s"
"    %s%s valC = %ssrcC%s"
"    %s%s destVal = valC ? valA : valB;\n"
"    %s"
"}\n";

static const char *kernel_qc_load_plain_prefix = "";
static const char *kernel_qc_load_plain_suffix = "[ tid ];\n";

static const char *kernel_qc_load_vec3_prefix = "vload3( tid, ";
static const char *kernel_qc_load_vec3_suffix = ");\n";

static const char *kernel_qc_store_plain = "dst[ tid ] = destVal;\n";
static const char *kernel_qc_store_vec3 = "vstore3( destVal, tid, dst );\n";

int test_question_colon_op(cl_device_id deviceID, cl_context context,
                           cl_command_queue queue, int num_elements,
                           int vectorSize, TestStyle style, ExplicitType type )
{
    cl_mem              streams[4];
    cl_int              *input_ptr[3], *output_ptr;
    cl_program          program;
    cl_kernel           kernel;
    size_t              threads[1];
    int                 err;
    int inputAVecSize, inputBVecSize, inputCVecSize;
    const char * sizeNames[] = { "", "", "2", "3", "4", "", "", "", "8", "", "", "", "", "", "", "", "16" };
    // Identical to sizeNames but with a blank for 3, since we use vload/store there
    const char * paramSizeNames[] = { "", "", "2", "", "4", "", "", "", "8", "", "", "", "", "", "", "", "16" };
    MTdata s_randStates;

    inputAVecSize = inputBVecSize = inputCVecSize = vectorSize;
    if( style & kInputCAlsoScalar )
    {
        style = (TestStyle)( style & ~kInputCAlsoScalar );
        inputCVecSize = 1;
    }
    if( style == kInputAScalar )
        inputAVecSize = 1;
    else if( style == kInputBScalar )
        inputBVecSize = 1;
    else if( style == kVectorScalarScalar )
        inputAVecSize = inputBVecSize = 1;

    log_info("Testing \"?:\" on %s%d (%s?%s:%s inputs)\n",
             get_explicit_type_name(type), vectorSize, ( inputCVecSize == 1 ) ? "scalar" : "vector",
             ( inputAVecSize == 1 ) ? "scalar" : "vector",
             ( inputBVecSize == 1 ) ? "scalar" : "vector" );


    const char *type_name = get_explicit_type_name(type);
    size_t type_size = get_explicit_type_size(type);

    // Create and initialize I/O buffers

    input_ptr[0] = (cl_int*)malloc(type_size * num_elements * vectorSize);
    input_ptr[1] = (cl_int*)malloc(type_size * num_elements * vectorSize);
    input_ptr[2] = (cl_int*)malloc(type_size * num_elements * vectorSize);
    output_ptr = (cl_int*)malloc(type_size * num_elements * vectorSize);

    s_randStates = init_genrand( gRandomSeed );

    generate_random_data( type, num_elements * inputAVecSize, s_randStates, input_ptr[ 0 ] );
    generate_random_data( type, num_elements * inputBVecSize, s_randStates, input_ptr[ 1 ] );
    generate_random_bool_data( num_elements * inputCVecSize, s_randStates, (cl_char *)input_ptr[ 2 ], type_size );

    streams[0] = clCreateBuffer(
        context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
        type_size * num_elements * inputAVecSize, input_ptr[0], &err);
    test_error(err, "clCreateBuffer failed");
    streams[1] = clCreateBuffer(
        context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
        type_size * num_elements * inputBVecSize, input_ptr[1], &err);
    test_error(err, "clCreateBuffer failed");
    streams[2] = clCreateBuffer(
        context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
        type_size * num_elements * inputCVecSize, input_ptr[2], &err);
    test_error(err, "clCreateBuffer failed");
    streams[3] =
        clCreateBuffer(context, CL_MEM_WRITE_ONLY,
                       type_size * num_elements * vectorSize, NULL, &err);
    test_error(err, "clCreateBuffer failed");

    const char *vectorString = sizeNames[ vectorSize ];
    const char *inputAVectorString = sizeNames[ inputAVecSize ];
    const char *inputBVectorString = sizeNames[ inputBVecSize ];
    const char *inputCVectorString = sizeNames[ inputCVecSize ];

    char programString[4096];
    const char *ptr;

    sprintf( programString, kernel_question_colon_full, type_name, paramSizeNames[ inputAVecSize ],
            type_name, paramSizeNames[ inputBVecSize ],
            type_name, paramSizeNames[ inputCVecSize ],
         type_name, paramSizeNames[ vectorSize ],
            // Loads
            type_name, inputAVectorString, ( inputAVecSize == 3 ) ? kernel_qc_load_vec3_prefix : kernel_qc_load_plain_prefix, ( inputAVecSize == 3 ) ? kernel_qc_load_vec3_suffix : kernel_qc_load_plain_suffix,
            type_name, inputBVectorString, ( inputBVecSize == 3 ) ? kernel_qc_load_vec3_prefix : kernel_qc_load_plain_prefix, ( inputBVecSize == 3 ) ? kernel_qc_load_vec3_suffix : kernel_qc_load_plain_suffix,
            type_name, inputCVectorString, ( inputCVecSize == 3 ) ? kernel_qc_load_vec3_prefix : kernel_qc_load_plain_prefix, ( inputCVecSize == 3 ) ? kernel_qc_load_vec3_suffix : kernel_qc_load_plain_suffix,
            // Dest type
            type_name, vectorString,
            // Store
            ( vectorSize == 3 ) ? kernel_qc_store_vec3 : kernel_qc_store_plain );

    ptr = programString;
    err = create_single_kernel_helper( context, &program, &kernel, 1, &ptr, "test" );
    test_error( err, "Unable to create test kernel" );

    err = clSetKernelArg( kernel, 0, sizeof streams[0], &streams[0] );
    err |= clSetKernelArg( kernel, 1, sizeof streams[1], &streams[1] );
    err |= clSetKernelArg( kernel, 2, sizeof streams[2], &streams[2] );
    err |= clSetKernelArg( kernel, 3, sizeof streams[3], &streams[3] );
    test_error(err, "clSetKernelArgs failed");

    // Run
    threads[0] = (size_t)num_elements;

    err = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
    test_error(err, "clEnqueueNDRangeKernel failed");

    // Read and verify results
    err = clEnqueueReadBuffer( queue, streams[3], CL_TRUE, 0, type_size*num_elements * vectorSize, (void *)output_ptr, 0, NULL, NULL );
    test_error(err, "clEnqueueReadBuffer failed");

    // log_info("Performing verification\n");
    int error_count = 0;

    char *inputAPtr = (char *)input_ptr[ 0 ];
    char *inputBPtr = (char *)input_ptr[ 1 ];
    cl_char *inputCPtr = (cl_char *)input_ptr[ 2 ];
    char *actualPtr = (char *)output_ptr;

    for( int i = 0; i < num_elements; i++ )
    {
        for( int j = 0; j < vectorSize; j++ )
        {
            char *expectedPtr = ( *inputCPtr ) ? inputAPtr : inputBPtr;
            if( memcmp( expectedPtr, actualPtr, type_size ) != 0 )
            {
#if 0
                char expectedStr[ 128 ], actualStr[ 128 ], inputAStr[ 128 ], inputBStr[ 128 ];
                print_type_to_string( type, inputAPtr, inputAStr );
                print_type_to_string( type, inputBPtr, inputBStr );
                print_type_to_string( type, expectedPtr, expectedStr );
                print_type_to_string( type, actualPtr, actualStr );
                log_error( "cl_%s verification failed at element %d:%d (expected %s, got %s, inputs: %s, %s, %s)\n",
                          type_name, i, j, expectedStr, actualStr, inputAStr, inputBStr, ( *inputCPtr ) ? "true" : "false" );
#endif
                error_count++;
            }
            // Advance for each element member. Note if any of the vec sizes are 1, they don't advance here
            inputAPtr += ( inputAVecSize == 1 ) ? 0 : type_size;
            inputBPtr += ( inputBVecSize == 1 ) ? 0 : type_size;
            inputCPtr += ( inputCVecSize == 1 ) ? 0 : type_size;
            actualPtr += ( vectorSize == 1 ) ? 0 : type_size;
        }
        // Reverse for the member advance. If the vec sizes are 1, we need to advance, but otherwise they're already correct
        inputAPtr += ( inputAVecSize == 1 ) ? type_size : 0;
        inputBPtr += ( inputBVecSize == 1 ) ? type_size : 0;
        inputCPtr += ( inputCVecSize == 1 ) ? type_size : 0;
        actualPtr += ( vectorSize == 1 ) ? type_size : 0;
    }

    // cleanup
    clReleaseMemObject(streams[0]);
    clReleaseMemObject(streams[1]);
    clReleaseMemObject(streams[2]);
    clReleaseMemObject(streams[3]);
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    free(input_ptr[0]);
    free(input_ptr[1]);
    free(input_ptr[2]);
    free(output_ptr);
    free_mtdata( s_randStates );

    return error_count;
}