//
// Copyright (c) 2017, 2020 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"

// Bug: Missing in spec: atomic_intptr_t is always supported if device is 32-bits.
// Bug: Missing in spec: CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE

#define FLUSH fflush(stdout)

#define MAX_STR 16*1024

#define ALIGNMENT 128

#define OPTIONS "-cl-std=CL2.0"

// NUM_ROUNDS must be at least 1.
// It determines how many sets of random data we push through the global
// variables.
#define NUM_ROUNDS 1

// This is a shared property of the writer and reader kernels.
#define NUM_TESTED_VALUES 5

// TODO: pointer-to-half (and its vectors)
// TODO: union of...

#include <algorithm>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <string>
#include <vector>
#include <cassert>
#include <sys/types.h>
#include <sys/stat.h>
#include "harness/typeWrappers.h"
#include "harness/errorHelpers.h"
#include "harness/mt19937.h"
#include "procs.h"


////////////////////
// Device capabilities
static int l_has_double = 0;
static int l_has_half = 0;
static int l_64bit_device = 0;
static int l_has_int64_atomics = 0;
static int l_has_intptr_atomics = 0;
static int l_has_cles_int64 = 0;

static int l_host_is_big_endian = 1;

static size_t l_max_global_id0 = 0;
static cl_bool l_linker_available = false;

#define check_error(errCode,msg,...) ((errCode != CL_SUCCESS) ? (log_error("ERROR: " msg "! (%s:%d)\n", ## __VA_ARGS__, __FILE__, __LINE__), 1) : 0)

////////////////////
// Info about types we can use for program scope variables.


class TypeInfo {

public:
    TypeInfo() :
        name(""),
        m_buf_elem_type(""),
        m_is_vecbase(false),
        m_is_atomic(false),
        m_is_like_size_t(false),
        m_is_bool(false),
        m_elem_type(0), m_num_elem(0),
        m_size(0),
        m_value_size(0)
        {}
    TypeInfo(const char* name_arg) :
        name(name_arg),
        m_buf_elem_type(name_arg),
        m_is_vecbase(false),
        m_is_atomic(false),
        m_is_like_size_t(false),
        m_is_bool(false),
        m_elem_type(0), m_num_elem(0),
        m_size(0),
        m_value_size(0)
        { }

    // Vectors
    TypeInfo( TypeInfo* elem_type, int num_elem ) :
        m_is_vecbase(false),
        m_is_atomic(false),
        m_is_like_size_t(false),
        m_is_bool(false),
        m_elem_type(elem_type),
        m_num_elem(num_elem)
    {
        char the_name[10]; // long enough for longest vector type name "double16"
        snprintf(the_name,sizeof(the_name),"%s%d",elem_type->get_name_c_str(),m_num_elem);
        this->name = std::string(the_name);
        this->m_buf_elem_type = std::string(the_name);
        this->m_value_size = num_elem * elem_type->get_size();
        if ( m_num_elem == 3 ) {
            this->m_size = 4 * elem_type->get_size();
        } else {
            this->m_size = num_elem * elem_type->get_size();
        }
    }
    const std::string& get_name(void) const { return name; }
    const char* get_name_c_str(void) const { return name.c_str(); }
    TypeInfo& set_vecbase(void) { this->m_is_vecbase = true; return *this; }
    TypeInfo& set_atomic(void) { this->m_is_atomic = true; return *this; }
    TypeInfo& set_like_size_t(void) {
        this->m_is_like_size_t = true;
        this->set_size( l_64bit_device ? 8 : 4 );
        this->m_buf_elem_type = l_64bit_device ? "ulong" : "uint";
        return *this;
    }
    TypeInfo& set_bool(void) { this->m_is_bool = true; return *this; }
    TypeInfo& set_size(size_t n) { this->m_value_size = this->m_size = n; return *this; }
    TypeInfo& set_buf_elem_type( const char* name ) { this->m_buf_elem_type = std::string(name); return *this; }

    const TypeInfo* elem_type(void) const { return m_elem_type; }
    int num_elem(void) const { return m_num_elem; }

    bool is_vecbase(void) const {return m_is_vecbase;}
    bool is_atomic(void) const {return m_is_atomic;}
    bool is_atomic_64bit(void) const {return m_is_atomic && m_size == 8;}
    bool is_like_size_t(void) const {return m_is_like_size_t;}
    bool is_bool(void) const {return m_is_bool;}
    size_t get_size(void) const {return m_size;}
    size_t get_value_size(void) const {return m_value_size;}

    // When passing values of this type to a kernel, what buffer type
    // should be used?
    const char* get_buf_elem_type(void) const { return m_buf_elem_type.c_str(); }

    std::string as_string(const cl_uchar* value_ptr) const {
        // This method would be shorter if I had a real handle to element
        // vector type.
        if ( this->is_bool() ) {
            std::string result( name );
            result += "<";
            result += (*value_ptr ? "true" : "false");
            result += ", ";
            char buf[10];
            sprintf(buf,"%02x",*value_ptr);
            result += buf;
            result += ">";
            return result;
        } else if ( this->num_elem() ) {
            std::string result( name );
            result += "<";
            for ( unsigned ielem = 0 ; ielem < this->num_elem() ; ielem++ ) {
                char buf[MAX_STR];
                if ( ielem ) result += ", ";
                for ( unsigned ibyte = 0; ibyte < this->m_elem_type->get_size() ; ibyte++ ) {
                    sprintf(buf + 2*ibyte,"%02x", value_ptr[ ielem * this->m_elem_type->get_size() + ibyte ] );
                }
                result += buf;
            }
            result += ">";
            return result;
        } else {
            std::string result( name );
            result += "<";
            char buf[MAX_STR];
            for ( unsigned ibyte = 0; ibyte < this->get_size() ; ibyte++ ) {
                sprintf(buf + 2*ibyte,"%02x", value_ptr[ ibyte ] );
            }
            result += buf;
            result += ">";
            return result;
        }
    }

    // Initialize the given buffer to a constant value initialized as if it
    // were from the INIT_VAR macro below.
    // Only needs to support values 0 and 1.
    void init( cl_uchar* buf, cl_uchar val) const {
        if ( this->num_elem() ) {
            for ( unsigned ielem = 0 ; ielem < this->num_elem() ; ielem++ ) {
                // Delegate!
                this->init_elem( buf + ielem * this->get_value_size()/this->num_elem(), val );
            }
        } else {
            init_elem( buf, val );
        }
    }

private:
    void init_elem( cl_uchar* buf, cl_uchar val ) const {
        size_t elem_size = this->num_elem() ? this->get_value_size()/this->num_elem() : this->get_size();
        memset(buf,0,elem_size);
        if ( val ) {
            if ( strstr( name.c_str(), "float" ) ) {
                *(float*)buf = (float)val;
                return;
            }
            if ( strstr( name.c_str(), "double" ) ) {
                *(double*)buf = (double)val;
                return;
            }
            if ( this->is_bool() ) { *buf = (bool)val; return; }

            // Write a single character value to the correct spot,
            // depending on host endianness.
            if ( l_host_is_big_endian ) *(buf + elem_size-1) = (cl_uchar)val;
            else *buf = (cl_uchar)val;
        }
    }
public:

    void dump(FILE* fp) const {
        fprintf(fp,"Type %s : <%d,%d,%s> ", name.c_str(),
                (int)m_size,
                (int)m_value_size,
                m_buf_elem_type.c_str() );
        if ( this->m_elem_type ) fprintf(fp, " vec(%s,%d)", this->m_elem_type->get_name_c_str(), this->num_elem() );
        if ( this->m_is_vecbase ) fprintf(fp, " vecbase");
        if ( this->m_is_bool ) fprintf(fp, " bool");
        if ( this->m_is_like_size_t ) fprintf(fp, " like-size_t");
        if ( this->m_is_atomic ) fprintf(fp, " atomic");
        fprintf(fp,"\n");
        fflush(fp);
    }

private:
    std::string name;
    TypeInfo* m_elem_type;
    int m_num_elem;
    bool m_is_vecbase;
    bool m_is_atomic;
    bool m_is_like_size_t;
    bool m_is_bool;
    size_t m_size; // Number of bytes of storage occupied by this type.
    size_t m_value_size; // Number of bytes of value significant for this type. Differs for vec3.

    // When passing values of this type to a kernel, what buffer type
    // should be used?
    // For most types, it's just itself.
    // Use a std::string so I don't have to make a copy constructor.
    std::string m_buf_elem_type;
};


#define NUM_SCALAR_TYPES (8+2) // signed and unsigned integral types, float and double
#define NUM_VECTOR_SIZES (5)   // 2,3,4,8,16
#define NUM_PLAIN_TYPES \
      5 /*boolean and size_t family */  \
    + NUM_SCALAR_TYPES \
    + NUM_SCALAR_TYPES*NUM_VECTOR_SIZES \
    + 10 /* atomic types */

// Need room for plain, array, pointer, struct
#define MAX_TYPES (4*NUM_PLAIN_TYPES)

static TypeInfo type_info[MAX_TYPES];
static int num_type_info = 0; // Number of valid entries in type_info[]




// A helper class to form kernel source arguments for clCreateProgramWithSource.
class StringTable {
public:
    StringTable() : m_c_strs(NULL), m_lengths(NULL), m_frozen(false), m_strings() {}
    ~StringTable() { release_frozen(); }

    void add(std::string s) { release_frozen(); m_strings.push_back(s); }

    const size_t num_str() { freeze(); return m_strings.size(); }
    const char** strs() { freeze(); return m_c_strs; }
    const size_t* lengths() { freeze(); return m_lengths; }

private:
    void freeze(void) {
        if ( !m_frozen ) {
            release_frozen();

            m_c_strs = (const char**) malloc(sizeof(const char*) * m_strings.size());
            m_lengths = (size_t*) malloc(sizeof(size_t) * m_strings.size());
            assert( m_c_strs );
            assert( m_lengths );

            for ( size_t i = 0; i < m_strings.size() ; i++ ) {
                m_c_strs[i] = m_strings[i].c_str();
                m_lengths[i] = strlen(m_c_strs[i]);
            }

            m_frozen = true;
        }
    }
    void release_frozen(void) {
        if ( m_c_strs ) { free(m_c_strs); m_c_strs = 0; }
        if ( m_lengths ) { free(m_lengths); m_lengths = 0; }
        m_frozen = false;
    }

    typedef std::vector<std::string> strlist_t;
    strlist_t m_strings;
    const char** m_c_strs;
    size_t* m_lengths;
    bool m_frozen;
};


////////////////////
// File scope function declarations

static void l_load_abilities(cl_device_id device);
static const char* l_get_fp64_pragma(void);
static const char* l_get_cles_int64_pragma(void);
static int l_build_type_table(cl_device_id device);

static int l_get_device_info(cl_device_id device, size_t* max_size_ret, size_t* pref_size_ret);

static void l_set_randomly( cl_uchar* buf, size_t buf_size, RandomSeed& rand_state );
static int l_compare( const cl_uchar* expected, const cl_uchar* received, unsigned num_values, const TypeInfo&ti );
static int l_copy( cl_uchar* dest, unsigned dest_idx, const cl_uchar* src, unsigned src_idx, const TypeInfo&ti );

static std::string conversion_functions(const TypeInfo& ti);
static std::string global_decls(const TypeInfo& ti, bool with_init);
static std::string global_check_function(const TypeInfo& ti);
static std::string writer_function(const TypeInfo& ti);
static std::string reader_function(const TypeInfo& ti);

static int l_write_read( cl_device_id device, cl_context context, cl_command_queue queue );
static int l_write_read_for_type( cl_device_id device, cl_context context, cl_command_queue queue, const TypeInfo& ti, RandomSeed& rand_state );

static int l_init_write_read( cl_device_id device, cl_context context, cl_command_queue queue );
static int l_init_write_read_for_type( cl_device_id device, cl_context context, cl_command_queue queue, const TypeInfo& ti, RandomSeed& rand_state );

static int l_capacity( cl_device_id device, cl_context context, cl_command_queue queue, size_t max_size );
static int l_user_type( cl_device_id device, cl_context context, cl_command_queue queue, size_t max_size, bool separate_compilation );



////////////////////
// File scope function definitions

static cl_int print_build_log(cl_program program, cl_uint num_devices, cl_device_id *device_list, cl_uint count, const char **strings, const size_t *lengths, const char* options)
{
    cl_uint i;
    cl_int error;
    BufferOwningPtr<cl_device_id> devices;

    if(num_devices == 0 || device_list == NULL)
    {
        error = clGetProgramInfo(program, CL_PROGRAM_NUM_DEVICES, sizeof(num_devices), &num_devices, NULL);
        test_error(error, "clGetProgramInfo CL_PROGRAM_NUM_DEVICES failed");

        device_list = (cl_device_id*)malloc(sizeof(cl_device_id)*num_devices);
        devices.reset(device_list);

        memset(device_list, 0, sizeof(cl_device_id) * num_devices);

        error = clGetProgramInfo(program, CL_PROGRAM_DEVICES, sizeof(cl_device_id) * num_devices, device_list, NULL);
        test_error(error, "clGetProgramInfo CL_PROGRAM_DEVICES failed");
    }

    cl_uint z;
    bool sourcePrinted = false;

    for(z = 0; z < num_devices; z++)
    {
        char deviceName[4096] = "";
        error = clGetDeviceInfo(device_list[z], CL_DEVICE_NAME, sizeof(deviceName), deviceName, NULL);
        check_error(error, "Device \"%d\" failed to return a name. clGetDeviceInfo CL_DEVICE_NAME failed", z);

        cl_build_status buildStatus;
        error = clGetProgramBuildInfo(program, device_list[z], CL_PROGRAM_BUILD_STATUS, sizeof(buildStatus), &buildStatus, NULL);
        check_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_STATUS failed");

        if(buildStatus != CL_BUILD_SUCCESS)
        {
            if(!sourcePrinted)
            {
                log_error("Build options: %s\n", options);
                if(count && strings)
                {
                    log_error("Original source is: ------------\n");
                    for(i = 0; i < count; i++) log_error("%s", strings[i]);
                }
                sourcePrinted = true;
            }

            char statusString[64] = "";
            if (buildStatus == (cl_build_status)CL_BUILD_SUCCESS)
              sprintf(statusString, "CL_BUILD_SUCCESS");
            else if (buildStatus == (cl_build_status)CL_BUILD_NONE)
              sprintf(statusString, "CL_BUILD_NONE");
            else if (buildStatus == (cl_build_status)CL_BUILD_ERROR)
              sprintf(statusString, "CL_BUILD_ERROR");
            else if (buildStatus == (cl_build_status)CL_BUILD_IN_PROGRESS)
              sprintf(statusString, "CL_BUILD_IN_PROGRESS");
            else
              sprintf(statusString, "UNKNOWN (%d)", buildStatus);

            log_error("Build not successful for device \"%s\", status: %s\n", deviceName, statusString);

            size_t paramSize = 0;
            error = clGetProgramBuildInfo(program, device_list[z], CL_PROGRAM_BUILD_LOG, 0, NULL, &paramSize);
            if(check_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_LOG failed")) break;

            std::string log;
            log.resize(paramSize/sizeof(char));

            error = clGetProgramBuildInfo(program, device_list[z], CL_PROGRAM_BUILD_LOG, paramSize, &log[0], NULL);
            if(check_error(error, "Device %d (%s) failed to return a build log", z, deviceName)) break;
            if(log[0] == 0) log_error("clGetProgramBuildInfo returned an empty log.\n");
            else
            {
                log_error("Build log:\n", deviceName);
                log_error("%s\n", log.c_str());
            }
        }
    }
    return error;
}

static void l_load_abilities(cl_device_id device)
{
    l_has_half       = is_extension_available(device,"cl_khr_fp16");
    l_has_double     = is_extension_available(device,"cl_khr_fp64");
    l_has_cles_int64 = is_extension_available(device,"cles_khr_int64");

    l_has_int64_atomics
    =  is_extension_available(device,"cl_khr_int64_base_atomics")
    && is_extension_available(device,"cl_khr_int64_extended_atomics");

    {
        int status = CL_SUCCESS;
        cl_uint addr_bits = 32;
        status = clGetDeviceInfo(device,CL_DEVICE_ADDRESS_BITS,sizeof(addr_bits),&addr_bits,0);
        l_64bit_device = ( status == CL_SUCCESS && addr_bits == 64 );
    }

    // 32-bit devices always have intptr atomics.
    l_has_intptr_atomics = !l_64bit_device || l_has_int64_atomics;

    union { char c[4]; int i; } probe;
    probe.i = 1;
    l_host_is_big_endian = !probe.c[0];

    // Determine max global id.
    {
        int status = CL_SUCCESS;
        cl_uint max_dim = 0;
        status = clGetDeviceInfo(device,CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,sizeof(max_dim),&max_dim,0);
        assert( status == CL_SUCCESS );
        assert( max_dim > 0 );
        size_t max_id[3];
        max_id[0] = 0;
    status = clGetDeviceInfo(device,CL_DEVICE_MAX_WORK_ITEM_SIZES,max_dim*sizeof(size_t),&max_id[0],0);
        assert( status == CL_SUCCESS );
        l_max_global_id0 = max_id[0];
    }

    { // Is separate compilation supported?
        int status = CL_SUCCESS;
        l_linker_available = false;
        status = clGetDeviceInfo(device,CL_DEVICE_LINKER_AVAILABLE,sizeof(l_linker_available),&l_linker_available,0);
        assert( status == CL_SUCCESS );
    }
}


static const char* l_get_fp64_pragma(void)
{
    return l_has_double ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" : "";
}

static const char* l_get_cles_int64_pragma(void)
{
    return l_has_cles_int64 ? "#pragma OPENCL EXTENSION cles_khr_int64 : enable\n" : "";
}

static const char* l_get_int64_atomic_pragma(void)
{
    return "#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable\n"
           "#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable\n";
}

static int l_build_type_table(cl_device_id device)
{
    int status = CL_SUCCESS;
    size_t iscalar = 0;
    size_t ivecsize = 0;
    int vecsizes[] = { 2, 3, 4, 8, 16 };
    const char* vecbase[] = {
        "uchar", "char",
        "ushort", "short",
        "uint", "int",
        "ulong", "long",
        "float",
        "double"
    };
    int vecbase_size[] = {
        1, 1,
        2, 2,
        4, 4,
        8, 8,
        4,
        8
    };
    const char* like_size_t[] = {
        "intptr_t",
        "uintptr_t",
        "size_t",
        "ptrdiff_t"
    };
    const char* atomics[] = {
        "atomic_int", "atomic_uint",
        "atomic_long", "atomic_ulong",
        "atomic_float",
        "atomic_double",
    };
    int atomics_size[] = {
        4, 4,
        8, 8,
        4,
        8
    };
    const char* intptr_atomics[] = {
        "atomic_intptr_t",
        "atomic_uintptr_t",
        "atomic_size_t",
        "atomic_ptrdiff_t"
    };

    l_load_abilities(device);
    num_type_info = 0;

    // Boolean.
    type_info[ num_type_info++ ] = TypeInfo( "bool" ).set_bool().set_size(1).set_buf_elem_type("uchar");

    // Vector types, and the related scalar element types.
    for ( iscalar=0; iscalar < sizeof(vecbase)/sizeof(vecbase[0]) ; ++iscalar ) {
        if ( !gHasLong && strstr(vecbase[iscalar],"long") ) continue;
        if ( !l_has_double && strstr(vecbase[iscalar],"double") ) continue;

        // Scalar
        TypeInfo* elem_type = type_info + num_type_info++;
        *elem_type = TypeInfo( vecbase[iscalar] ).set_vecbase().set_size( vecbase_size[iscalar] );

        // Vector
        for ( ivecsize=0; ivecsize < sizeof(vecsizes)/sizeof(vecsizes[0]) ; ivecsize++ ) {
            type_info[ num_type_info++ ] = TypeInfo( elem_type, vecsizes[ivecsize] );
        }
    }

    // Size_t-like types
    for ( iscalar=0; iscalar < sizeof(like_size_t)/sizeof(like_size_t[0]) ; ++iscalar ) {
        type_info[ num_type_info++ ] = TypeInfo( like_size_t[iscalar] ).set_like_size_t();
    }

    // Atomic types.
    for ( iscalar=0; iscalar < sizeof(atomics)/sizeof(atomics[0]) ; ++iscalar ) {
        if ( !l_has_int64_atomics && strstr(atomics[iscalar],"long") ) continue;
        if ( !(l_has_int64_atomics && l_has_double) && strstr(atomics[iscalar],"double") ) continue;

        // The +7 is used to skip over the "atomic_" prefix.
        const char* buf_type = atomics[iscalar] + 7;
        type_info[ num_type_info++ ] = TypeInfo( atomics[iscalar] ).set_atomic().set_size( atomics_size[iscalar] ).set_buf_elem_type( buf_type );
    }
    if ( l_has_intptr_atomics ) {
        for ( iscalar=0; iscalar < sizeof(intptr_atomics)/sizeof(intptr_atomics[0]) ; ++iscalar ) {
            type_info[ num_type_info++ ] = TypeInfo( intptr_atomics[iscalar] ).set_atomic().set_like_size_t();
        }
    }

    assert( num_type_info <= MAX_TYPES ); // or increase MAX_TYPES

#if 0
    for ( size_t i = 0 ; i < num_type_info ; i++ ) {
        type_info[ i ].dump(stdout);
    }
    exit(0);
#endif

    return status;
}

static const TypeInfo& l_find_type( const char* name )
{
    auto itr =
        std::find_if(type_info, type_info + num_type_info,
                     [name](TypeInfo& ti) { return ti.get_name() == name; });
    assert(itr != type_info + num_type_info);
    return *itr;
}



// Populate return parameters for max program variable size, preferred program variable size.

static int l_get_device_info(cl_device_id device, size_t* max_size_ret, size_t* pref_size_ret)
{
    int err = CL_SUCCESS;
    size_t return_size = 0;

    err = clGetDeviceInfo(device, CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE, sizeof(*max_size_ret), max_size_ret, &return_size);
    if ( err != CL_SUCCESS ) {
        log_error("Error: Failed to get device info for CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE\n");
        return err;
    }
    if ( return_size != sizeof(size_t) ) {
        log_error("Error: Invalid size %d returned for CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE\n", (int)return_size );
        return 1;
    }
    if ( return_size != sizeof(size_t) ) {
        log_error("Error: Invalid size %d returned for CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE\n", (int)return_size );
        return 1;
    }

    return_size = 0;
    err = clGetDeviceInfo(device, CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE, sizeof(*pref_size_ret), pref_size_ret, &return_size);
    if ( err != CL_SUCCESS ) {
        log_error("Error: Failed to get device info for CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE: %d\n",err);
        return err;
    }
    if ( return_size != sizeof(size_t) ) {
        log_error("Error: Invalid size %d returned for CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE\n", (int)return_size );
        return 1;
    }

    return CL_SUCCESS;
}


static void l_set_randomly( cl_uchar* buf, size_t buf_size, RandomSeed& rand_state )
{
    assert( 0 == (buf_size % sizeof(cl_uint) ) );
    for ( size_t i = 0; i < buf_size ; i += sizeof(cl_uint) ) {
        *( (cl_uint*)(buf + i) ) = genrand_int32( rand_state );
    }
#if 0
    for ( size_t i = 0; i < buf_size ; i++ ) {
        printf("%02x",buf[i]);
    }
    printf("\n");
#endif
}

// Return num_value values of the given type.
// Returns CL_SUCCESS if they compared as equal.
static int l_compare( const char* test_name, const cl_uchar* expected, const cl_uchar* received, size_t num_values, const TypeInfo&ti )
{
    // Compare only the valid returned bytes.
    for ( unsigned value_idx = 0; value_idx < num_values; value_idx++ ) {
        const cl_uchar* expv = expected + value_idx * ti.get_size();
        const cl_uchar* gotv = received + value_idx * ti.get_size();
        if ( memcmp( expv, gotv, ti.get_value_size() ) ) {
            std::string exp_str = ti.as_string( expv );
            std::string got_str = ti.as_string( gotv );
            log_error("Error: %s test for type %s, at index %d: Expected %s got %s\n",
                    test_name,
                    ti.get_name_c_str(), value_idx,
                    exp_str.c_str(),
                    got_str.c_str() );
            return 1;
        }
    }
    return CL_SUCCESS;
}

// Copy a target value from src[idx] to dest[idx]
static int l_copy( cl_uchar* dest, unsigned dest_idx, const cl_uchar* src, unsigned src_idx, const TypeInfo&ti )
{
    cl_uchar* raw_dest      = dest + dest_idx * ti.get_size();
    const cl_uchar* raw_src =  src +  src_idx * ti.get_size();
    memcpy( raw_dest, raw_src,  ti.get_value_size() );

    return 0;
}


static std::string conversion_functions(const TypeInfo& ti)
{
    std::string result;
    static char buf[MAX_STR];
    int num_printed = 0;
    // The atomic types just use the base type.
    if ( ti.is_atomic() || 0 == strcmp( ti.get_buf_elem_type(), ti.get_name_c_str() ) ) {
        // The type is represented in a buffer by itself.
        num_printed = snprintf(buf,MAX_STR,
                "%s from_buf(%s a) { return a; }\n"
                "%s to_buf(%s a) { return a; }\n",
                ti.get_buf_elem_type(), ti.get_buf_elem_type(),
                ti.get_buf_elem_type(), ti.get_buf_elem_type() );
    } else {
        // Just use C-style cast.
        num_printed = snprintf(buf,MAX_STR,
                "%s from_buf(%s a) { return (%s)a; }\n"
                "%s to_buf(%s a) { return (%s)a; }\n",
                ti.get_name_c_str(), ti.get_buf_elem_type(), ti.get_name_c_str(),
                ti.get_buf_elem_type(), ti.get_name_c_str(), ti.get_buf_elem_type() );
    }
    // Add initializations.
    if ( ti.is_atomic() ) {
        num_printed += snprintf( buf + num_printed, MAX_STR-num_printed,
                "#define INIT_VAR(a) ATOMIC_VAR_INIT(a)\n" );
    } else {
        // This cast works even if the target type is a vector type.
        num_printed += snprintf( buf + num_printed, MAX_STR-num_printed,
                "#define INIT_VAR(a) ((%s)(a))\n", ti.get_name_c_str());
    }
    assert( num_printed < MAX_STR ); // or increase MAX_STR
    result = buf;
    return result;
}

static std::string global_decls(const TypeInfo& ti, bool with_init )
{
    const char* tn = ti.get_name_c_str();
    const char* vol = (ti.is_atomic() ? " volatile " : " ");
    static char decls[MAX_STR];
    int num_printed = 0;
    if ( with_init ) {
        const char *decls_template_with_init =
            "%s %s var = INIT_VAR(0);\n"
            "global %s %s g_var = INIT_VAR(1);\n"
            "%s %s a_var[2] = { INIT_VAR(1), INIT_VAR(1) };\n"
            "volatile global %s %s* p_var = &a_var[1];\n\n";
        num_printed = snprintf(decls,sizeof(decls),decls_template_with_init,
                vol,tn,vol,tn,vol,tn,vol,tn);
    } else {
        const char *decls_template_no_init =
            "%s %s var;\n"
            "global %s %s g_var;\n"
            "%s %s a_var[2];\n"
            "global %s %s* p_var;\n\n";
        num_printed = snprintf(decls,sizeof(decls),decls_template_no_init,
             vol,tn,vol,tn,vol,tn,vol,tn);
    }
    assert( num_printed < sizeof(decls) );
    return std::string(decls);
}

// Return the source code for the "global_check" function for the given type.
// This function checks that all program-scope variables have appropriate
// initial values when no explicit initializer is used. If all tests pass the
// kernel writes a non-zero value to its output argument, otherwise it writes
// zero.
static std::string global_check_function(const TypeInfo& ti)
{
    const std::string type_name = ti.get_buf_elem_type();

    // all() should only be used on vector inputs. For scalar comparison, the
    // result of the equality operator can be used as a bool value.
    const bool is_scalar = ti.num_elem() == 0; // 0 is used to represent scalar types, not 1.
    const std::string is_equality_true = is_scalar ? "" : "all";

    std::string code = "kernel void global_check(global int* out) {\n";
    code += "  const " + type_name + " zero = ((" + type_name + ")0);\n";
    code += "  bool status = true;\n";
    if (ti.is_atomic()) {
        code += "  status &= " + is_equality_true + "(atomic_load(&var) == zero);\n";
        code += "  status &= " + is_equality_true + "(atomic_load(&g_var) == zero);\n";
        code += "  status &= " + is_equality_true + "(atomic_load(&a_var[0]) == zero);\n";
        code += "  status &= " + is_equality_true + "(atomic_load(&a_var[1]) == zero);\n";
    } else {
        code += "  status &= " + is_equality_true + "(var == zero);\n";
        code += "  status &= " + is_equality_true + "(g_var == zero);\n";
        code += "  status &= " + is_equality_true + "(a_var[0] == zero);\n";
        code += "  status &= " + is_equality_true + "(a_var[1] == zero);\n";
    }
    code += "  status &= (p_var == NULL);\n";
    code += "  *out = status ? 1 : 0;\n";
    code += "}\n\n";

    return code;
}

// Return the source text for the writer function for the given type.
// For types that can't be passed as pointer-to-type as a kernel argument,
// use a substitute base type of the same size.
static std::string writer_function(const TypeInfo& ti)
{
    static char writer_src[MAX_STR];
    int num_printed = 0;
    if ( !ti.is_atomic() ) {
        const char* writer_template_normal =
            "kernel void writer( global %s* src, uint idx ) {\n"
            "  var = from_buf(src[0]);\n"
            "  g_var = from_buf(src[1]);\n"
            "  a_var[0] = from_buf(src[2]);\n"
            "  a_var[1] = from_buf(src[3]);\n"
            "  p_var = a_var + idx;\n"
            "}\n\n";
        num_printed = snprintf(writer_src,sizeof(writer_src),writer_template_normal,ti.get_buf_elem_type());
    } else {
        const char* writer_template_atomic =
            "kernel void writer( global %s* src, uint idx ) {\n"
            "  atomic_store( &var, from_buf(src[0]) );\n"
            "  atomic_store( &g_var, from_buf(src[1]) );\n"
            "  atomic_store( &a_var[0], from_buf(src[2]) );\n"
            "  atomic_store( &a_var[1], from_buf(src[3]) );\n"
            "  p_var = a_var + idx;\n"
            "}\n\n";
        num_printed = snprintf(writer_src,sizeof(writer_src),writer_template_atomic,ti.get_buf_elem_type());
    }
    assert( num_printed < sizeof(writer_src) );
    std::string result = writer_src;
    return result;
}


// Return source text for teh reader function for the given type.
// For types that can't be passed as pointer-to-type as a kernel argument,
// use a substitute base type of the same size.
static std::string reader_function(const TypeInfo& ti)
{
    static char reader_src[MAX_STR];
    int num_printed = 0;
    if ( !ti.is_atomic() ) {
        const char* reader_template_normal =
            "kernel void reader( global %s* dest, %s ptr_write_val ) {\n"
            "  *p_var = from_buf(ptr_write_val);\n"
            "  dest[0] = to_buf(var);\n"
            "  dest[1] = to_buf(g_var);\n"
            "  dest[2] = to_buf(a_var[0]);\n"
            "  dest[3] = to_buf(a_var[1]);\n"
            "}\n\n";
        num_printed = snprintf(reader_src,sizeof(reader_src),reader_template_normal,ti.get_buf_elem_type(),ti.get_buf_elem_type());
    } else {
        const char* reader_template_atomic =
            "kernel void reader( global %s* dest, %s ptr_write_val ) {\n"
            "  atomic_store( p_var, from_buf(ptr_write_val) );\n"
            "  dest[0] = to_buf( atomic_load( &var ) );\n"
            "  dest[1] = to_buf( atomic_load( &g_var ) );\n"
            "  dest[2] = to_buf( atomic_load( &a_var[0] ) );\n"
            "  dest[3] = to_buf( atomic_load( &a_var[1] ) );\n"
            "}\n\n";
        num_printed = snprintf(reader_src,sizeof(reader_src),reader_template_atomic,ti.get_buf_elem_type(),ti.get_buf_elem_type());
    }
    assert( num_printed < sizeof(reader_src) );
    std::string result = reader_src;
    return result;
}

// Check that all globals where appropriately default-initialized.
static int check_global_initialization(cl_context context, cl_program program, cl_command_queue queue)
{
    int status = CL_SUCCESS;

    // Create a buffer on device to store a unique integer.
    cl_int is_init_valid = 0;
    clMemWrapper buffer(clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(is_init_valid), &is_init_valid, &status));
    test_error_ret(status, "Failed to allocate buffer", status);

    // Create, setup and invoke kernel.
    clKernelWrapper global_check(clCreateKernel(program, "global_check", &status));
    test_error_ret(status, "Failed to create global_check kernel", status);
    status = clSetKernelArg(global_check, 0, sizeof(cl_mem), &buffer);
    test_error_ret(status, "Failed to set up argument for the global_check kernel", status);
    const cl_uint work_dim = 1;
    const size_t global_work_offset[] = { 0 };
    const size_t global_work_size[] = { 1 };
    status = clEnqueueNDRangeKernel(queue, global_check, work_dim, global_work_offset, global_work_size, nullptr, 0, nullptr, nullptr);
    test_error_ret(status, "Failed to run global_check kernel", status);
    status = clFinish(queue);
    test_error_ret(status, "clFinish() failed", status);

    // Read back the memory buffer from the device.
    status = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, sizeof(is_init_valid), &is_init_valid, 0, nullptr, nullptr);
    test_error_ret(status, "Failed to read buffer from device", status);
    if (is_init_valid == 0) {
        log_error("Unexpected default values were detected");
        return 1;
    }

    return CL_SUCCESS;
}

// Check write-then-read.
static int l_write_read( cl_device_id device, cl_context context, cl_command_queue queue )
{
    int status = CL_SUCCESS;
    int itype;

    RandomSeed rand_state( gRandomSeed );

    for ( itype = 0; itype < num_type_info ; itype++ ) {
        status = status | l_write_read_for_type(device,context,queue,type_info[itype], rand_state );
        FLUSH;
    }

    return status;
}

static int l_write_read_for_type( cl_device_id device, cl_context context, cl_command_queue queue, const TypeInfo& ti, RandomSeed& rand_state )
{
    int err = CL_SUCCESS;
    std::string type_name( ti.get_name() );
    const char* tn = type_name.c_str();
    log_info("  %s ",tn);

    StringTable ksrc;
    ksrc.add( l_get_fp64_pragma() );
    ksrc.add( l_get_cles_int64_pragma() );
    if (ti.is_atomic_64bit())
      ksrc.add( l_get_int64_atomic_pragma() );
    ksrc.add( conversion_functions(ti) );
    ksrc.add( global_decls(ti,false) );
    ksrc.add( global_check_function(ti) );
    ksrc.add( writer_function(ti) );
    ksrc.add( reader_function(ti) );

    int status = CL_SUCCESS;
    clProgramWrapper program;
    clKernelWrapper writer;

    status = create_single_kernel_helper_with_build_options(context, &program, &writer, ksrc.num_str(), ksrc.strs(), "writer", OPTIONS);
    test_error_ret(status,"Failed to create program for read-after-write test",status);

    clKernelWrapper reader( clCreateKernel( program, "reader", &status ) );
    test_error_ret(status,"Failed to create reader kernel for read-after-write test",status);

    // Check size query.
    size_t used_bytes = 0;
    status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
    test_error_ret(status,"Failed to query global variable total size",status);
    size_t expected_used_bytes =
        (NUM_TESTED_VALUES-1)*ti.get_size() // Two regular variables and an array of 2 elements.
        + ( l_64bit_device ? 8 : 4 ); // The pointer
    if ( used_bytes < expected_used_bytes ) {
        log_error("Error program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)expected_used_bytes, (unsigned long long)used_bytes );
        err |= 1;
    }

    err |= check_global_initialization(context, program, queue);

    // We need to create 5 random values of the given type,
    // and read 4 of them back.
    const size_t write_data_size = NUM_TESTED_VALUES * sizeof(cl_ulong16);
    const size_t read_data_size = (NUM_TESTED_VALUES - 1) * sizeof(cl_ulong16);
    cl_uchar* write_data = (cl_uchar*)align_malloc(write_data_size, ALIGNMENT);
    cl_uchar* read_data = (cl_uchar*)align_malloc(read_data_size, ALIGNMENT);

    clMemWrapper write_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, write_data_size, write_data, &status ) );
    test_error_ret(status,"Failed to allocate write buffer",status);
    clMemWrapper read_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, read_data_size, read_data, &status ) );
    test_error_ret(status,"Failed to allocate read buffer",status);

    status = clSetKernelArg(writer,0,sizeof(cl_mem),&write_mem); test_error_ret(status,"set arg",status);
    status = clSetKernelArg(reader,0,sizeof(cl_mem),&read_mem); test_error_ret(status,"set arg",status);

    // Boolean random data needs to be massaged a bit more.
    const int num_rounds = ti.is_bool() ? (1 << NUM_TESTED_VALUES ) : NUM_ROUNDS;
    unsigned bool_iter = 0;

    for ( int iround = 0; iround < num_rounds ; iround++ ) {
        for ( cl_uint iptr_idx = 0; iptr_idx < 2 ; iptr_idx++ ) { // Index into array, to write via pointer
            // Generate new random data to push through.
            // Generate 5 * 128 bytes all the time, even though the test for many types use less than all that.

            cl_uchar *write_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, write_mem, CL_TRUE, CL_MAP_WRITE, 0, write_data_size, 0, 0, 0, 0);

            if ( ti.is_bool() ) {
                // For boolean, random data cast to bool isn't very random.
                // So use the bottom bit of bool_value_iter to get true
                // diversity.
                for ( unsigned value_idx = 0; value_idx < NUM_TESTED_VALUES ; value_idx++ ) {
                    write_data[value_idx] = (1<<value_idx) & bool_iter;
                    //printf(" %s", (write_data[value_idx] ? "true" : "false" ));
                }
                bool_iter++;
            } else {
                l_set_randomly( write_data, write_data_size, rand_state );
            }
            status = clSetKernelArg(writer,1,sizeof(cl_uint),&iptr_idx); test_error_ret(status,"set arg",status);

            // The value to write via the pointer should be taken from the
            // 5th typed slot of the write_data.
            status = clSetKernelArg(reader,1,ti.get_size(),write_data + (NUM_TESTED_VALUES-1)*ti.get_size()); test_error_ret(status,"set arg",status);

            // Determine the expected values.
            cl_uchar expected[read_data_size];
            memset( expected, -1, sizeof(expected) );
            l_copy( expected, 0, write_data, 0, ti );
            l_copy( expected, 1, write_data, 1, ti );
            l_copy( expected, 2, write_data, 2, ti );
            l_copy( expected, 3, write_data, 3, ti );
            // But we need to take into account the value from the pointer write.
            // The 2 represents where the "a" array values begin in our read-back.
            l_copy( expected, 2 + iptr_idx, write_data, 4, ti );

            clEnqueueUnmapMemObject(queue, write_mem, write_ptr, 0, 0, 0);

            if ( ti.is_bool() ) {
                // Collapse down to one bit.
                for ( unsigned i = 0; i <  NUM_TESTED_VALUES-1 ; i++ ) expected[i] = (bool)expected[i];
            }

            cl_uchar *read_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, read_mem, CL_TRUE, CL_MAP_READ, 0, read_data_size, 0, 0, 0, 0);
            memset(read_data, -1, read_data_size);
            clEnqueueUnmapMemObject(queue, read_mem, read_ptr, 0, 0, 0);

            // Now run the kernel
            const size_t one = 1;
            status = clEnqueueNDRangeKernel(queue,writer,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue writer",status);
            status = clEnqueueNDRangeKernel(queue,reader,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue reader",status);
            status = clFinish(queue); test_error_ret(status,"finish",status);

            read_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, read_mem, CL_TRUE, CL_MAP_READ, 0, read_data_size, 0, 0, 0, 0);

            if ( ti.is_bool() ) {
                // Collapse down to one bit.
                for ( unsigned i = 0; i <  NUM_TESTED_VALUES-1 ; i++ ) read_data[i] = (bool)read_data[i];
            }

            // Compare only the valid returned bytes.
            int compare_result = l_compare( "read-after-write", expected, read_data, NUM_TESTED_VALUES-1, ti );
            // log_info("Compared %d values each of size %llu. Result %d\n", NUM_TESTED_VALUES-1, (unsigned long long)ti.get_value_size(), compare_result );
            err |= compare_result;

            clEnqueueUnmapMemObject(queue, read_mem, read_ptr, 0, 0, 0);

            if ( err ) break;
        }
    }

    if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }
    align_free(write_data);
    align_free(read_data);
    return err;
}


// Check initialization, then, read, then write, then read.
static int l_init_write_read( cl_device_id device, cl_context context, cl_command_queue queue )
{
    int status = CL_SUCCESS;
    int itype;

    RandomSeed rand_state( gRandomSeed );

    for ( itype = 0; itype < num_type_info ; itype++ ) {
        status = status | l_init_write_read_for_type(device,context,queue,type_info[itype], rand_state );
    }
    return status;
}
static int l_init_write_read_for_type( cl_device_id device, cl_context context, cl_command_queue queue, const TypeInfo& ti, RandomSeed& rand_state )
{
    int err = CL_SUCCESS;
    std::string type_name( ti.get_name() );
    const char* tn = type_name.c_str();
    log_info("  %s ",tn);

    StringTable ksrc;
    ksrc.add( l_get_fp64_pragma() );
    ksrc.add( l_get_cles_int64_pragma() );
    if (ti.is_atomic_64bit())
      ksrc.add( l_get_int64_atomic_pragma() );
    ksrc.add( conversion_functions(ti) );
    ksrc.add( global_decls(ti,true) );
    ksrc.add( writer_function(ti) );
    ksrc.add( reader_function(ti) );

    int status = CL_SUCCESS;
    clProgramWrapper program;
    clKernelWrapper writer;

    status = create_single_kernel_helper_with_build_options(context, &program, &writer, ksrc.num_str(), ksrc.strs(), "writer", OPTIONS);
    test_error_ret(status,"Failed to create program for init-read-after-write test",status);

    clKernelWrapper reader( clCreateKernel( program, "reader", &status ) );
    test_error_ret(status,"Failed to create reader kernel for init-read-after-write test",status);

    // Check size query.
    size_t used_bytes = 0;
    status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
    test_error_ret(status,"Failed to query global variable total size",status);
    size_t expected_used_bytes =
        (NUM_TESTED_VALUES-1)*ti.get_size() // Two regular variables and an array of 2 elements.
        + ( l_64bit_device ? 8 : 4 ); // The pointer
    if ( used_bytes < expected_used_bytes ) {
        log_error("Error: program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)expected_used_bytes, (unsigned long long)used_bytes );
        err |= 1;
    }

    // We need to create 5 random values of the given type,
    // and read 4 of them back.
    const size_t write_data_size = NUM_TESTED_VALUES * sizeof(cl_ulong16);
    const size_t read_data_size = (NUM_TESTED_VALUES-1) * sizeof(cl_ulong16);

    cl_uchar* write_data = (cl_uchar*)align_malloc(write_data_size, ALIGNMENT);
    cl_uchar* read_data = (cl_uchar*)align_malloc(read_data_size, ALIGNMENT);
    clMemWrapper write_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, write_data_size, write_data, &status ) );
    test_error_ret(status,"Failed to allocate write buffer",status);
    clMemWrapper read_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, read_data_size, read_data, &status ) );
    test_error_ret(status,"Failed to allocate read buffer",status);

    status = clSetKernelArg(writer,0,sizeof(cl_mem),&write_mem); test_error_ret(status,"set arg",status);
    status = clSetKernelArg(reader,0,sizeof(cl_mem),&read_mem); test_error_ret(status,"set arg",status);

    // Boolean random data needs to be massaged a bit more.
    const int num_rounds = ti.is_bool() ? (1 << NUM_TESTED_VALUES ) : NUM_ROUNDS;
    unsigned bool_iter = 0;

    // We need to count iterations.  We do something *different on the
    // first iteration, to ensure we actually pick up the initialized
    // values.
    unsigned iteration = 0;

    for ( int iround = 0; iround < num_rounds ; iround++ ) {
        for ( cl_uint iptr_idx = 0; iptr_idx < 2 ; iptr_idx++ ) { // Index into array, to write via pointer
            // Generate new random data to push through.
            // Generate 5 * 128 bytes all the time, even though the test for many types use less than all that.

            cl_uchar *write_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, write_mem, CL_TRUE, CL_MAP_WRITE, 0, write_data_size, 0, 0, 0, 0);

            if ( ti.is_bool() ) {
                // For boolean, random data cast to bool isn't very random.
                // So use the bottom bit of bool_value_iter to get true
                // diversity.
                for ( unsigned value_idx = 0; value_idx < NUM_TESTED_VALUES ; value_idx++ ) {
                    write_data[value_idx] = (1<<value_idx) & bool_iter;
                    //printf(" %s", (write_data[value_idx] ? "true" : "false" ));
                }
                bool_iter++;
            } else {
                l_set_randomly( write_data, write_data_size, rand_state );
            }
            status = clSetKernelArg(writer,1,sizeof(cl_uint),&iptr_idx); test_error_ret(status,"set arg",status);

            if ( !iteration ) {
                // On first iteration, the value we write via the last arg
                // to the "reader" function is 0.
                // It's way easier to code the test this way.
                ti.init( write_data + (NUM_TESTED_VALUES-1)*ti.get_size(), 0 );
            }

            // The value to write via the pointer should be taken from the
            // 5th typed slot of the write_data.
            status = clSetKernelArg(reader,1,ti.get_size(),write_data + (NUM_TESTED_VALUES-1)*ti.get_size()); test_error_ret(status,"set arg",status);

            // Determine the expected values.
            cl_uchar expected[read_data_size];
            memset( expected, -1, sizeof(expected) );
            if ( iteration ) {
                l_copy( expected, 0, write_data, 0, ti );
                l_copy( expected, 1, write_data, 1, ti );
                l_copy( expected, 2, write_data, 2, ti );
                l_copy( expected, 3, write_data, 3, ti );
                // But we need to take into account the value from the pointer write.
                // The 2 represents where the "a" array values begin in our read-back.
                // But we need to take into account the value from the pointer write.
                l_copy( expected, 2 + iptr_idx, write_data, 4, ti );
            } else {
                // On first iteration, expect these initialized values!
                // See the decls_template_with_init above.
                ti.init( expected, 0 );
                ti.init( expected + ti.get_size(), 1 );
                ti.init( expected + 2*ti.get_size(), 1 );
                // Emulate the effect of the write via the pointer.
                // The value is 0, not 1 (see above).
                // The pointer is always initialized to the second element
                // of the array. So it goes into slot 3 of the "expected" array.
                ti.init( expected + 3*ti.get_size(), 0 );
            }

            if ( ti.is_bool() ) {
                // Collapse down to one bit.
                for ( unsigned i = 0; i <  NUM_TESTED_VALUES-1 ; i++ ) expected[i] = (bool)expected[i];
            }

            clEnqueueUnmapMemObject(queue, write_mem, write_ptr, 0, 0, 0);

            cl_uchar *read_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, read_mem, CL_TRUE, CL_MAP_READ, 0, read_data_size, 0, 0, 0, 0);
            memset( read_data, -1, read_data_size );
            clEnqueueUnmapMemObject(queue, read_mem, read_ptr, 0, 0, 0);

            // Now run the kernel
            const size_t one = 1;
            if ( iteration ) {
                status = clEnqueueNDRangeKernel(queue,writer,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue writer",status);
            } else {
                // On first iteration, we should be picking up the
                // initialized value. So don't enqueue the writer.
            }
            status = clEnqueueNDRangeKernel(queue,reader,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue reader",status);
            status = clFinish(queue); test_error_ret(status,"finish",status);

            read_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, read_mem, CL_TRUE, CL_MAP_READ, 0, read_data_size, 0, 0, 0, 0);

            if ( ti.is_bool() ) {
                // Collapse down to one bit.
                for ( unsigned i = 0; i <  NUM_TESTED_VALUES-1 ; i++ ) read_data[i] = (bool)read_data[i];
            }

            // Compare only the valid returned bytes.
            //log_info(" Round %d ptr_idx %u\n", iround, iptr_idx );
            int compare_result = l_compare( "init-write-read", expected, read_data, NUM_TESTED_VALUES-1, ti );
            //log_info("Compared %d values each of size %llu. Result %d\n", NUM_TESTED_VALUES-1, (unsigned long long)ti.get_value_size(), compare_result );
            err |= compare_result;

            clEnqueueUnmapMemObject(queue, read_mem, read_ptr, 0, 0, 0);

            if ( err ) break;

            iteration++;
        }
    }

    if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }
    align_free(write_data);
    align_free(read_data);

    return err;
}


// Check that we can make at least one variable with size
// max_size which is returned from the device info property : CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE.
static int l_capacity( cl_device_id device, cl_context context, cl_command_queue queue, size_t max_size )
{
    int err = CL_SUCCESS;
    // Just test one type.
    const TypeInfo ti( l_find_type("uchar") );
    log_info(" l_capacity...");

    const char prog_src_template[] =
#if defined(_WIN32)
        "uchar var[%Iu];\n\n"
#else
        "uchar var[%zu];\n\n"
#endif
        "kernel void get_max_size( global ulong* size_ret ) {\n"
#if defined(_WIN32)
        "  *size_ret = (ulong)%Iu;\n"
#else
        "  *size_ret = (ulong)%zu;\n"
#endif
        "}\n\n"
        "kernel void writer( global uchar* src ) {\n"
        "  var[get_global_id(0)] = src[get_global_linear_id()];\n"
        "}\n\n"
        "kernel void reader( global uchar* dest ) {\n"
        "  dest[get_global_linear_id()] = var[get_global_id(0)];\n"
        "}\n\n";
    char prog_src[MAX_STR];
    int num_printed = snprintf(prog_src,sizeof(prog_src),prog_src_template,max_size, max_size);
    assert( num_printed < MAX_STR ); // or increase MAX_STR

    StringTable ksrc;
    ksrc.add( prog_src );

    int status = CL_SUCCESS;
    clProgramWrapper program;
    clKernelWrapper get_max_size;

    status = create_single_kernel_helper_with_build_options(context, &program, &get_max_size, ksrc.num_str(), ksrc.strs(), "get_max_size", OPTIONS);
    test_error_ret(status,"Failed to create program for capacity test",status);

    // Check size query.
    size_t used_bytes = 0;
    status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
    test_error_ret(status,"Failed to query global variable total size",status);
    if ( used_bytes < max_size ) {
        log_error("Error: program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)max_size, (unsigned long long)used_bytes );
        err |= 1;
    }

    // Prepare to execute
    clKernelWrapper writer( clCreateKernel( program, "writer", &status ) );
    test_error_ret(status,"Failed to create writer kernel for capacity test",status);
    clKernelWrapper reader( clCreateKernel( program, "reader", &status ) );
    test_error_ret(status,"Failed to create reader kernel for capacity test",status);

    cl_ulong max_size_ret = 0;
    const size_t arr_size = 10*1024*1024;
    cl_uchar* buffer = (cl_uchar*) align_malloc( arr_size, ALIGNMENT );

    if ( !buffer ) { log_error("Failed to allocate buffer\n"); return 1; }

    clMemWrapper max_size_ret_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, sizeof(max_size_ret), &max_size_ret, &status ) );
    test_error_ret(status,"Failed to allocate size query buffer",status);
    clMemWrapper buffer_mem( clCreateBuffer( context, CL_MEM_READ_WRITE, arr_size, 0, &status ) );
    test_error_ret(status,"Failed to allocate write buffer",status);

    status = clSetKernelArg(get_max_size,0,sizeof(cl_mem),&max_size_ret_mem); test_error_ret(status,"set arg",status);
    status = clSetKernelArg(writer,0,sizeof(cl_mem),&buffer_mem); test_error_ret(status,"set arg",status);
    status = clSetKernelArg(reader,0,sizeof(cl_mem),&buffer_mem); test_error_ret(status,"set arg",status);

    // Check the macro value of CL_DEVICE_MAX_GLOBAL_VARIABLE
    const size_t one = 1;
    status = clEnqueueNDRangeKernel(queue,get_max_size,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue size query",status);
    status = clFinish(queue); test_error_ret(status,"finish",status);

    cl_uchar *max_size_ret_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, max_size_ret_mem, CL_TRUE, CL_MAP_READ, 0, sizeof(max_size_ret), 0, 0, 0, 0);
    if ( max_size_ret != max_size ) {
        log_error("Error: preprocessor definition for CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE is %llu and does not match device query value %llu\n",
                (unsigned long long) max_size_ret,
                (unsigned long long) max_size );
        err |= 1;
    }
    clEnqueueUnmapMemObject(queue, max_size_ret_mem, max_size_ret_ptr, 0, 0, 0);

    RandomSeed rand_state_write( gRandomSeed );
    for ( size_t offset = 0; offset < max_size ; offset += arr_size ) {
        size_t curr_size = (max_size - offset) < arr_size ? (max_size - offset) : arr_size;
        l_set_randomly( buffer, curr_size, rand_state_write );
        status = clEnqueueWriteBuffer (queue, buffer_mem, CL_TRUE, 0, curr_size, buffer, 0, 0, 0);test_error_ret(status,"populate buffer_mem object",status);
        status = clEnqueueNDRangeKernel(queue,writer,1,&offset,&curr_size,0,0,0,0); test_error_ret(status,"enqueue writer",status);
    status = clFinish(queue); test_error_ret(status,"finish",status);
    }

    RandomSeed rand_state_read( gRandomSeed );
    for ( size_t offset = 0; offset < max_size ; offset += arr_size ) {
        size_t curr_size = (max_size - offset) < arr_size ? (max_size - offset) : arr_size;
        status = clEnqueueNDRangeKernel(queue,reader,1,&offset,&curr_size,0,0,0,0); test_error_ret(status,"enqueue reader",status);
        cl_uchar* read_mem_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, buffer_mem, CL_TRUE, CL_MAP_READ, 0, curr_size, 0, 0, 0, &status);test_error_ret(status,"map read data",status);
        l_set_randomly( buffer, curr_size, rand_state_read );
        err |= l_compare( "capacity", buffer, read_mem_ptr, curr_size, ti );
        clEnqueueUnmapMemObject(queue, buffer_mem, read_mem_ptr, 0, 0, 0);
    }

    if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }
    align_free(buffer);

    return err;
}


// Check operation on a user type.
static int l_user_type( cl_device_id device, cl_context context, cl_command_queue queue, bool separate_compile )
{
    int err = CL_SUCCESS;
    // Just test one type.
    const TypeInfo ti( l_find_type("uchar") );
    log_info(" l_user_type %s...", separate_compile ? "separate compilation" : "single source compilation" );

    if ( separate_compile && ! l_linker_available ) {
        log_info("Separate compilation is not supported. Skipping test\n");
        return err;
    }

    const char type_src[] =
        "typedef struct { uchar c; uint i; } my_struct_t;\n\n";
    const char def_src[] =
        "my_struct_t var = { 'a', 42 };\n\n";
    const char decl_src[] =
        "extern my_struct_t var;\n\n";

    // Don't use a host struct. We can't guarantee that the host
    // compiler has the same structure layout as the device compiler.
    const char writer_src[] =
        "kernel void writer( uchar c, uint i ) {\n"
        "  var.c = c;\n"
        "  var.i = i;\n"
        "}\n\n";
    const char reader_src[] =
        "kernel void reader( global uchar* C, global uint* I ) {\n"
        "  *C = var.c;\n"
        "  *I = var.i;\n"
        "}\n\n";

    clProgramWrapper program;

    if ( separate_compile ) {
        // Separate compilation flow.
        StringTable wksrc;
        wksrc.add( type_src );
        wksrc.add( def_src );
        wksrc.add( writer_src );

        StringTable rksrc;
        rksrc.add( type_src );
        rksrc.add( decl_src );
        rksrc.add( reader_src );

        int status = CL_SUCCESS;
        clProgramWrapper writer_program( clCreateProgramWithSource( context, wksrc.num_str(), wksrc.strs(), wksrc.lengths(), &status ) );
        test_error_ret(status,"Failed to create writer program for user type test",status);

        status = clCompileProgram( writer_program, 1, &device, OPTIONS, 0, 0, 0, 0, 0 );
        if(check_error(status, "Failed to compile writer program for user type test (%s)", IGetErrorString(status)))
        {
            print_build_log(writer_program, 1, &device, wksrc.num_str(), wksrc.strs(), wksrc.lengths(), OPTIONS);
            return status;
        }

        clProgramWrapper reader_program( clCreateProgramWithSource( context, rksrc.num_str(), rksrc.strs(), rksrc.lengths(), &status ) );
        test_error_ret(status,"Failed to create reader program for user type test",status);

        status = clCompileProgram( reader_program, 1, &device, OPTIONS, 0, 0, 0, 0, 0 );
        if(check_error(status, "Failed to compile reader program for user type test (%s)", IGetErrorString(status)))
        {
            print_build_log(reader_program, 1, &device, rksrc.num_str(), rksrc.strs(), rksrc.lengths(), OPTIONS);
            return status;
        }

        cl_program progs[2];
        progs[0] = writer_program;
        progs[1] = reader_program;

        program = clLinkProgram( context, 1, &device, "", 2, progs, 0, 0, &status );
        if(check_error(status, "Failed to link program for user type test (%s)", IGetErrorString(status)))
        {
            print_build_log(program, 1, &device, 0, NULL, NULL, "");
            return status;
        }
    } else {
        // Single compilation flow.
        StringTable ksrc;
        ksrc.add( type_src );
        ksrc.add( def_src );
        ksrc.add( writer_src );
        ksrc.add( reader_src );

        int status = CL_SUCCESS;

        status = create_single_kernel_helper_create_program(context, &program, ksrc.num_str(), ksrc.strs(), OPTIONS);
        if(check_error(status, "Failed to build program for user type test (%s)", IGetErrorString(status)))
        {
            print_build_log(program, 1, &device, ksrc.num_str(), ksrc.strs(), ksrc.lengths(), OPTIONS);
            return status;
        }

        status = clBuildProgram(program, 1, &device, OPTIONS, 0, 0);
        if(check_error(status, "Failed to compile program for user type test (%s)", IGetErrorString(status)))
        {
            print_build_log(program, 1, &device, ksrc.num_str(), ksrc.strs(), ksrc.lengths(), OPTIONS);
            return status;
        }
    }


    // Check size query.
    size_t used_bytes = 0;
    int status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
    test_error_ret(status,"Failed to query global variable total size",status);
    size_t expected_size = sizeof(cl_uchar) + sizeof(cl_uint);
    if ( used_bytes < expected_size ) {
        log_error("Error: program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)expected_size, (unsigned long long)used_bytes );
        err |= 1;
    }

    // Prepare to execute
    clKernelWrapper writer( clCreateKernel( program, "writer", &status ) );
    test_error_ret(status,"Failed to create writer kernel for user type test",status);
    clKernelWrapper reader( clCreateKernel( program, "reader", &status ) );
    test_error_ret(status,"Failed to create reader kernel for user type test",status);

    // Set up data.
    cl_uchar* uchar_data = (cl_uchar*)align_malloc(sizeof(cl_uchar), ALIGNMENT);
    cl_uint* uint_data = (cl_uint*)align_malloc(sizeof(cl_uint), ALIGNMENT);

    clMemWrapper uchar_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, sizeof(cl_uchar), uchar_data, &status ) );
    test_error_ret(status,"Failed to allocate uchar buffer",status);
    clMemWrapper uint_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, sizeof(cl_uint), uint_data, &status ) );
    test_error_ret(status,"Failed to allocate uint buffer",status);

    status = clSetKernelArg(reader,0,sizeof(cl_mem),&uchar_mem); test_error_ret(status,"set arg",status);
    status = clSetKernelArg(reader,1,sizeof(cl_mem),&uint_mem); test_error_ret(status,"set arg",status);

    cl_uchar expected_uchar = 'a';
    cl_uint expected_uint = 42;
    for ( unsigned iter = 0; iter < 5 ; iter++ ) { // Must go around at least twice
        // Read back data
        *uchar_data = -1;
        *uint_data = -1;
        const size_t one = 1;
        status = clEnqueueNDRangeKernel(queue,reader,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue reader",status);
        status = clFinish(queue); test_error_ret(status,"finish",status);

        cl_uchar *uint_data_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, uint_mem, CL_TRUE, CL_MAP_READ, 0, sizeof(cl_uint), 0, 0, 0, 0);
        cl_uchar *uchar_data_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, uchar_mem, CL_TRUE, CL_MAP_READ, 0, sizeof(cl_uchar), 0, 0, 0, 0);

        if ( expected_uchar != *uchar_data || expected_uint != *uint_data ) {
            log_error("FAILED: Iteration %d Got (0x%2x,%d) but expected (0x%2x,%d)\n",
                    iter, (int)*uchar_data, *uint_data, (int)expected_uchar, expected_uint );
            err |= 1;
        }

        clEnqueueUnmapMemObject(queue, uint_mem, uint_data_ptr, 0, 0, 0);
        clEnqueueUnmapMemObject(queue, uchar_mem, uchar_data_ptr, 0, 0, 0);

        // Mutate the data.
        expected_uchar++;
        expected_uint++;

        // Write the new values into persistent store.
        *uchar_data = expected_uchar;
        *uint_data = expected_uint;
        status = clSetKernelArg(writer,0,sizeof(cl_uchar),uchar_data); test_error_ret(status,"set arg",status);
        status = clSetKernelArg(writer,1,sizeof(cl_uint),uint_data); test_error_ret(status,"set arg",status);
        status = clEnqueueNDRangeKernel(queue,writer,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue writer",status);
        status = clFinish(queue); test_error_ret(status,"finish",status);
    }

    if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }
    align_free(uchar_data);
    align_free(uint_data);
    return err;
}

// Determines whether its valid to skip this test based on the driver version
// and the features it optionally supports.
// Whether the test should be skipped is writen into the out paramter skip.
// The check returns an error code for the clDeviceInfo query.
static cl_int should_skip(cl_device_id device, cl_bool& skip)
{
    // Assume we can't skip to begin with.
    skip = CL_FALSE;

    // Progvar tests are already skipped for OpenCL < 2.0, so here we only need
    // to test for 3.0 since that is when program scope global variables become
    // optional.
    if (get_device_cl_version(device) >= Version(3, 0))
    {
        size_t max_global_variable_size{};
        test_error(clGetDeviceInfo(device, CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE,
                                   sizeof(max_global_variable_size),
                                   &max_global_variable_size, nullptr),
                   "clGetDeviceInfo failed");
        skip = (max_global_variable_size != 0) ? CL_FALSE : CL_TRUE;
    }
    return CL_SUCCESS;
}

////////////////////
// Global functions


// Test support for variables at program scope. Miscellaneous
int test_progvar_prog_scope_misc(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
    cl_bool skip{ CL_FALSE };
    auto error = should_skip(device, skip);
    if (CL_SUCCESS != error)
    {
        return TEST_FAIL;
    }
    if (skip)
    {
        log_info("Skipping progvar_prog_scope_misc since it is optionally not "
                 "supported on this device\n");
        return TEST_SKIPPED_ITSELF;
    }
    size_t max_size = 0;
    size_t pref_size = 0;

    cl_int err = CL_SUCCESS;

    err = l_get_device_info( device, &max_size, &pref_size );
    err |= l_build_type_table( device );

    err |= l_capacity( device, context, queue, max_size );
    err |= l_user_type( device, context, queue, false );
    err |= l_user_type( device, context, queue, true );

    return err;
}


// Test support for variables at program scope. Unitialized data
int test_progvar_prog_scope_uninit(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
    cl_bool skip{ CL_FALSE };
    auto error = should_skip(device, skip);
    if (CL_SUCCESS != error)
    {
        return TEST_FAIL;
    }
    if (skip)
    {
        log_info(
            "Skipping progvar_prog_scope_uninit since it is optionally not "
            "supported on this device\n");
        return TEST_SKIPPED_ITSELF;
    }
    size_t max_size = 0;
    size_t pref_size = 0;

    cl_int err = CL_SUCCESS;

    err = l_get_device_info( device, &max_size, &pref_size );
    err |= l_build_type_table( device );

    err |= l_write_read( device, context, queue );

    return err;
}

// Test support for variables at program scope. Initialized data.
int test_progvar_prog_scope_init(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
    cl_bool skip{ CL_FALSE };
    auto error = should_skip(device, skip);
    if (CL_SUCCESS != error)
    {
        return TEST_FAIL;
    }
    if (skip)
    {
        log_info("Skipping progvar_prog_scope_init since it is optionally not "
                 "supported on this device\n");
        return TEST_SKIPPED_ITSELF;
    }
    size_t max_size = 0;
    size_t pref_size = 0;

    cl_int err = CL_SUCCESS;

    err = l_get_device_info( device, &max_size, &pref_size );
    err |= l_build_type_table( device );

    err |= l_init_write_read( device, context, queue );

    return err;
}


// A simple test for support of static variables inside a kernel.
int test_progvar_func_scope(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
    cl_bool skip{ CL_FALSE };
    auto error = should_skip(device, skip);
    if (CL_SUCCESS != error)
    {
        return TEST_FAIL;
    }
    if (skip)
    {
        log_info("Skipping progvar_func_scope since it is optionally not "
                 "supported on this device\n");
        return TEST_SKIPPED_ITSELF;
    }
    size_t max_size = 0;
    size_t pref_size = 0;

    cl_int err = CL_SUCCESS;

    // Deliberately have two variables with the same name but in different
    // scopes.
    // Also, use a large initialized structure in both cases.
    const char prog_src[] =
        "typedef struct { char c; int16 i; } mystruct_t;\n"
        "kernel void test_bump( global int* value, int which ) {\n"
        "  if ( which ) {\n"
        // Explicit address space.
        // Last element set to 0
        "     static global mystruct_t persistent = {'a',(int16)(0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,0) };\n"
        "     *value = persistent.i.sf++;\n"
        "  } else {\n"
        // Implicitly global
        // Last element set to 100
        "     static mystruct_t persistent = {'b',(int16)(0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,100) };\n"
        "     *value = persistent.i.sf++;\n"
        "  }\n"
        "}\n";

    StringTable ksrc;
    ksrc.add( prog_src );

    int status = CL_SUCCESS;
    clProgramWrapper program;
    clKernelWrapper test_bump;

    status = create_single_kernel_helper_with_build_options(context, &program, &test_bump, ksrc.num_str(), ksrc.strs(), "test_bump", OPTIONS);
    test_error_ret(status, "Failed to create program for function static variable test", status);

    // Check size query.
    size_t used_bytes = 0;
    status = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE, sizeof(used_bytes), &used_bytes, 0 );
    test_error_ret(status,"Failed to query global variable total size",status);
    size_t expected_size = 2 * sizeof(cl_int); // Two ints.
    if ( used_bytes < expected_size ) {
        log_error("Error: program query for global variable total size query failed: Expected at least %llu but got %llu\n", (unsigned long long)expected_size, (unsigned long long)used_bytes );
        err |= 1;
    }

    // Prepare the data.
    cl_int counter_value = 0;
    clMemWrapper counter_value_mem( clCreateBuffer( context, CL_MEM_USE_HOST_PTR, sizeof(counter_value), &counter_value, &status ) );
    test_error_ret(status,"Failed to allocate counter query buffer",status);

    status = clSetKernelArg(test_bump,0,sizeof(cl_mem),&counter_value_mem); test_error_ret(status,"set arg",status);

    // Go a few rounds, alternating between the two counters in the kernel.

    // Same as initial values in kernel.
    // But "true" which increments the 0-based counter, and "false" which
    // increments the 100-based counter.
    cl_int expected_counter[2] = { 100, 0 };

    const size_t one = 1;
    for ( int iround = 0; iround < 5 ; iround++ ) { // Must go at least twice around
        for ( int iwhich = 0; iwhich < 2 ; iwhich++ ) { // Cover both counters
            status = clSetKernelArg(test_bump,1,sizeof(iwhich),&iwhich); test_error_ret(status,"set arg",status);
            status = clEnqueueNDRangeKernel(queue,test_bump,1,0,&one,0,0,0,0); test_error_ret(status,"enqueue test_bump",status);
            status = clFinish(queue); test_error_ret(status,"finish",status);

            cl_uchar *counter_value_ptr = (cl_uchar *)clEnqueueMapBuffer(queue, counter_value_mem, CL_TRUE, CL_MAP_READ, 0, sizeof(counter_value), 0, 0, 0, 0);

            if ( counter_value != expected_counter[iwhich] ) {
                log_error("Error: Round %d on counter %d: Expected %d but got %d\n",
                        iround, iwhich, expected_counter[iwhich], counter_value );
                err |= 1;
            }
            expected_counter[iwhich]++; // Emulate behaviour of the kernel.

            clEnqueueUnmapMemObject(queue, counter_value_mem, counter_value_ptr, 0, 0, 0);
        }
    }

    if ( CL_SUCCESS == err ) { log_info("OK\n"); FLUSH; }

    return err;
}