OpenCL-CTS/test_conformance/gl/test_images_read_common.cpp

//
// Copyright (c) 2017 The Khronos Group Inc.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
//    http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include "common.h"
#include "testBase.h"

#if defined(__APPLE__)
#include <OpenGL/glu.h>
#else
#include <GL/glu.h>
#include <CL/cl_gl.h>
#endif

extern int supportsHalf(cl_context context, bool *supports_half);
extern int supportsMsaa(cl_context context, bool *supports_msaa);
extern int supportsDepth(cl_context context, bool *supports_depth);

// clang-format off
static const char *kernelpattern_image_read_1d =
"__kernel void sample_test( read_only image1d_t source, sampler_t sampler, __global %s4 *results )\n"
"{\n"
"  int offset = get_global_id(0);\n"
"  results[ offset ] = read_image%s( source, sampler, offset );\n"
"}\n";

static const char *kernelpattern_image_read_1d_buffer =
"__kernel void sample_test( read_only image1d_buffer_t source, sampler_t sampler, __global %s4 *results )\n"
"{\n"
"  int offset = get_global_id(0);\n"
"  results[ offset ] = read_image%s( source, offset );\n"
"}\n";

static const char *kernelpattern_image_read_1darray =
"__kernel void sample_test( read_only image1d_array_t source, sampler_t sampler, __global %s4 *results )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    results[ tidY * get_image_width( source ) + tidX ] = read_image%s( source, sampler, (int2)( tidX, tidY ) );\n"
"}\n";

static const char *kernelpattern_image_read_2d =
"__kernel void sample_test( read_only image2d_t source, sampler_t sampler, __global %s4 *results )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    results[ tidY * get_image_width( source ) + tidX ] = read_image%s( source, sampler, (int2)( tidX, tidY ) );\n"
"}\n";

static const char *kernelpattern_image_read_2darray =
"__kernel void sample_test( read_only image2d_array_t source, sampler_t sampler, __global %s4 *results )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    int  tidZ = get_global_id(2);\n"
"    int  width = get_image_width( source );\n"
"    int  height = get_image_height( source );\n"
"    int offset = tidZ * width * height + tidY * width + tidX;\n"
"\n"
"     results[ offset ] = read_image%s( source, sampler, (int4)( tidX, tidY, tidZ, 0 ) );\n"
"}\n";

static const char *kernelpattern_image_read_3d =
"__kernel void sample_test( read_only image3d_t source, sampler_t sampler, __global %s4 *results )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    int  tidZ = get_global_id(2);\n"
"    int  width = get_image_width( source );\n"
"    int  height = get_image_height( source );\n"
"    int offset = tidZ * width * height + tidY * width + tidX;\n"
"\n"
"     results[ offset ] = read_image%s( source, sampler, (int4)( tidX, tidY, tidZ, 0 ) );\n"
"}\n";

static const char *kernelpattern_image_read_2d_depth =
"__kernel void sample_test( read_only image2d_depth_t source, sampler_t sampler, __global %s *results )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    results[ tidY * get_image_width( source ) + tidX ] = read_image%s( source, sampler, (int2)( tidX, tidY ) );\n"
"}\n";

static const char *kernelpattern_image_read_2darray_depth =
"__kernel void sample_test( read_only image2d_array_depth_t source, sampler_t sampler, __global %s *results )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    int  tidZ = get_global_id(2);\n"
"    int  width = get_image_width( source );\n"
"    int  height = get_image_height( source );\n"
"    int offset = tidZ * width * height + tidY * width + tidX;\n"
"\n"
"     results[ offset ] = read_image%s( source, sampler, (int4)( tidX, tidY, tidZ, 0 ) );\n"
"}\n";

static const char *kernelpattern_image_multisample_read_2d =
"#pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable\n"
"__kernel void sample_test( read_only image2d_msaa_t source, sampler_t sampler, __global %s4 *results )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    int  width = get_image_width( source );\n"
"    int  height = get_image_height( source );\n"
"    int  num_samples = get_image_num_samples( source );\n"
"    for(size_t sample = 0; sample < num_samples; sample++ ) {\n"
"    int  offset = sample * width * height + tidY * width + tidX;\n"
"     results[ offset ] = read_image%s( source, (int2)( tidX, tidY ), sample );\n"
"    }\n"
"}\n";

static const char *kernelpattern_image_multisample_read_2d_depth =
  "#pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable\n"
  "__kernel void sample_test( read_only image2d_msaa_depth_t source, sampler_t sampler, __global %s *results )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    int  width = get_image_width( source );\n"
"    int  height = get_image_height( source );\n"
  "    int  num_samples = get_image_num_samples( source );\n"
  "    for(size_t sample = 0; sample < num_samples; sample++ ) {\n"
"    int  offset = sample * width * height + tidY * width + tidX;\n"
"     results[ offset ] = read_image%s( source, (int2)( tidX, tidY ), sample );\n"
  "    }\n"
"}\n";

static const char *kernelpattern_image_multisample_read_2darray =
"#pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable\n"
"__kernel void sample_test( read_only image2d_array_msaa_t source, sampler_t sampler, __global %s4 *results )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    int  tidZ = get_global_id(2);\n"
"    int  num_samples = get_image_num_samples( source );\n"
"    int  width  = get_image_width( source );\n"
"    int  height = get_image_height( source );\n"
"    int  array_size = get_image_array_size( source );\n"
"    for(size_t sample = 0; sample< num_samples; ++sample) {\n"
"      int offset = (array_size * width * height) * sample + (width * height) * tidZ + tidY * width + tidX;\n"
"         results[ offset ] = read_image%s( source, (int4)( tidX, tidY, tidZ, 1 ), sample );\n"
"    }\n"
"}\n";

static const char *kernelpattern_image_multisample_read_2darray_depth =
  "#pragma OPENCL EXTENSION cl_khr_gl_msaa_sharing : enable\n"
  "__kernel void sample_test( read_only image2d_array_msaa_depth_t source, sampler_t sampler, __global %s *results )\n"
"{\n"
"    int  tidX = get_global_id(0);\n"
"    int  tidY = get_global_id(1);\n"
"    int  tidZ = get_global_id(2);\n"
"    int  num_samples = get_image_num_samples( source );\n"
"    int  width  = get_image_width( source );\n"
"    int  height = get_image_height( source );\n"
  "    int  array_size = get_image_array_size( source );\n"
  "    for(size_t sample = 0; sample < num_samples; ++sample) {\n"
  "      int offset = (array_size * width * height) * sample + (width * height) * tidZ + tidY * width + tidX;\n"
  "         results[ offset ] = read_image%s( source, (int4)( tidX, tidY, tidZ, 1 ), sample );\n"
  "    }\n"
"}\n";
// clang-format on

static const char *
get_appropriate_kernel_for_target(GLenum target, cl_channel_order channel_order)
{

    switch (get_base_gl_target(target))
    {
        case GL_TEXTURE_1D: return kernelpattern_image_read_1d;
        case GL_TEXTURE_BUFFER: return kernelpattern_image_read_1d_buffer;
        case GL_TEXTURE_1D_ARRAY: return kernelpattern_image_read_1darray;
        case GL_TEXTURE_RECTANGLE_EXT:
        case GL_TEXTURE_2D:
        case GL_COLOR_ATTACHMENT0:
        case GL_RENDERBUFFER:
        case GL_TEXTURE_CUBE_MAP:
#ifdef GL_VERSION_3_2
            if (channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
                return kernelpattern_image_read_2d_depth;
#endif
            return kernelpattern_image_read_2d;
        case GL_TEXTURE_2D_ARRAY:
#ifdef GL_VERSION_3_2
            if (channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
                return kernelpattern_image_read_2darray_depth;
#endif
            return kernelpattern_image_read_2darray;
        case GL_TEXTURE_3D: return kernelpattern_image_read_3d;
        case GL_TEXTURE_2D_MULTISAMPLE:
#ifdef GL_VERSION_3_2
            if (channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
                return kernelpattern_image_multisample_read_2d_depth;
#endif
            return kernelpattern_image_multisample_read_2d;
            break;
        case GL_TEXTURE_2D_MULTISAMPLE_ARRAY:
#ifdef GL_VERSION_3_2
            if (channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
                return kernelpattern_image_multisample_read_2darray_depth;
#endif
            return kernelpattern_image_multisample_read_2darray;
            break;
        default:
            log_error("Unsupported texture target (%s); cannot determine "
                      "appropriate kernel.",
                      GetGLTargetName(target));
            return NULL;
    }
}

int test_cl_image_read(cl_context context, cl_command_queue queue,
                       GLenum gl_target, cl_mem image, size_t width,
                       size_t height, size_t depth, size_t sampleNum,
                       cl_image_format *outFormat, ExplicitType *outType,
                       void **outResultBuffer)
{
    clProgramWrapper program;
    clKernelWrapper kernel;
    clMemWrapper streams[2];

    int error;
    char kernelSource[2048];
    char *programPtr;

    // Use the image created from the GL texture.
    streams[0] = image;

    // Determine data type and format that CL came up with
    error = clGetImageInfo(streams[0], CL_IMAGE_FORMAT, sizeof(cl_image_format),
                           outFormat, NULL);
    test_error(error, "Unable to get CL image format");

    // Determine the number of samples
    cl_uint samples = 0;
    error = clGetImageInfo(streams[0], CL_IMAGE_NUM_SAMPLES, sizeof(samples),
                           &samples, NULL);
    test_error(error, "Unable to get CL_IMAGE_NUM_SAMPLES");

    // Create the source
    *outType = get_read_kernel_type(outFormat);
    size_t channelSize = get_explicit_type_size(*outType);

    const char *source = get_appropriate_kernel_for_target(
        gl_target, outFormat->image_channel_order);

    sprintf(kernelSource, source, get_explicit_type_name(*outType),
            get_kernel_suffix(outFormat));

    programPtr = kernelSource;
    if (create_single_kernel_helper(context, &program, &kernel, 1,
                                    (const char **)&programPtr, "sample_test",
                                    ""))
    {
        return -1;
    }

    // Create a vanilla output buffer
    cl_device_id device;
    error = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, sizeof(device),
                                  &device, NULL);
    test_error(error, "Unable to get queue device");

    cl_ulong maxAllocSize = 0;
    error = clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
                            sizeof(maxAllocSize), &maxAllocSize, NULL);
    test_error(error, "Unable to get CL_DEVICE_MAX_MEM_ALLOC_SIZE");

    size_t buffer_bytes = channelSize
        * get_channel_order_channel_count(outFormat->image_channel_order)
        * width * height * depth * sampleNum;
    if (buffer_bytes > maxAllocSize)
    {
        log_info("Output buffer size %d is too large for device (max alloc "
                 "size %d) Skipping...\n",
                 (int)buffer_bytes, (int)maxAllocSize);
        return 1;
    }

    streams[1] =
        clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_bytes, NULL, &error);
    test_error(error, "Unable to create output buffer");

    /* Assign streams and execute */
    clSamplerWrapper sampler = clCreateSampler(
        context, CL_FALSE, CL_ADDRESS_NONE, CL_FILTER_NEAREST, &error);
    test_error(error, "Unable to create sampler");

    error = clSetKernelArg(kernel, 0, sizeof(streams[0]), &streams[0]);
    test_error(error, "Unable to set kernel arguments");
    error = clSetKernelArg(kernel, 1, sizeof(sampler), &sampler);
    test_error(error, "Unable to set kernel arguments");
    error = clSetKernelArg(kernel, 2, sizeof(streams[1]), &streams[1]);
    test_error(error, "Unable to set kernel arguments");

    glFinish();

    error =
        (*clEnqueueAcquireGLObjects_ptr)(queue, 1, &streams[0], 0, NULL, NULL);
    test_error(error, "Unable to acquire GL obejcts");

    // The ND range we use is a function of the dimensionality of the image.
    size_t global_range[3] = { width, height, depth };
    size_t *local_range = NULL;
    int ndim = 1;

    switch (get_base_gl_target(gl_target))
    {
        case GL_TEXTURE_1D:
        case GL_TEXTURE_BUFFER: ndim = 1; break;
        case GL_TEXTURE_RECTANGLE_EXT:
        case GL_TEXTURE_2D:
        case GL_TEXTURE_1D_ARRAY:
        case GL_COLOR_ATTACHMENT0:
        case GL_RENDERBUFFER:
        case GL_TEXTURE_CUBE_MAP: ndim = 2; break;
        case GL_TEXTURE_3D:
        case GL_TEXTURE_2D_ARRAY:
#ifdef GL_VERSION_3_2
        case GL_TEXTURE_2D_MULTISAMPLE:
        case GL_TEXTURE_2D_MULTISAMPLE_ARRAY: ndim = 3; break;
#endif
        default:
            log_error("Test error: Unsupported texture target.\n");
            return 1;
    }

    // 2D and 3D images have a special way to set the local size (legacy).
    // Otherwise, we let CL select by leaving local_range as NULL.

    if (gl_target == GL_TEXTURE_2D)
    {
        local_range = (size_t *)malloc(sizeof(size_t) * ndim);
        get_max_common_2D_work_group_size(context, kernel, global_range,
                                          local_range);
    }
    else if (gl_target == GL_TEXTURE_3D)
    {
        local_range = (size_t *)malloc(sizeof(size_t) * ndim);
        get_max_common_3D_work_group_size(context, kernel, global_range,
                                          local_range);
    }

    error = clEnqueueNDRangeKernel(queue, kernel, ndim, NULL, global_range,
                                   local_range, 0, NULL, NULL);
    test_error(error, "Unable to execute test kernel");

    error =
        (*clEnqueueReleaseGLObjects_ptr)(queue, 1, &streams[0], 0, NULL, NULL);
    test_error(error, "clEnqueueReleaseGLObjects failed");

    // Read results from the CL buffer
    *outResultBuffer = (void *)(new char[channelSize
                                         * get_channel_order_channel_count(
                                             outFormat->image_channel_order)
                                         * width * height * depth * sampleNum]);
    error = clEnqueueReadBuffer(
        queue, streams[1], CL_TRUE, 0,
        channelSize
            * get_channel_order_channel_count(outFormat->image_channel_order)
            * width * height * depth * sampleNum,
        *outResultBuffer, 0, NULL, NULL);
    test_error(error, "Unable to read output CL buffer!");

    // free the ranges
    if (local_range) free(local_range);

    return 0;
}

static int test_image_read(cl_context context, cl_command_queue queue,
                           GLenum target, GLuint globj, size_t width,
                           size_t height, size_t depth, size_t sampleNum,
                           cl_image_format *outFormat, ExplicitType *outType,
                           void **outResultBuffer)
{
    int error;

    // Create a CL image from the supplied GL texture or renderbuffer.
    cl_mem image;
    if (target == GL_RENDERBUFFER || target == GL_COLOR_ATTACHMENT0)
    {
        image = (*clCreateFromGLRenderbuffer_ptr)(context, CL_MEM_READ_ONLY,
                                                  globj, &error);
    }
    else
    {
        image = (*clCreateFromGLTexture_ptr)(context, CL_MEM_READ_ONLY, target,
                                             0, globj, &error);
    }

    if (error != CL_SUCCESS)
    {
        if (target == GL_RENDERBUFFER || target == GL_COLOR_ATTACHMENT0)
        {
            print_error(error,
                        "Unable to create CL image from GL renderbuffer");
        }
        else
        {
            print_error(error, "Unable to create CL image from GL texture");
            GLint fmt;
            glGetTexLevelParameteriv(target, 0, GL_TEXTURE_INTERNAL_FORMAT,
                                     &fmt);
            log_error("    Supplied GL texture was base format %s and internal "
                      "format %s\n",
                      GetGLBaseFormatName(fmt), GetGLFormatName(fmt));
        }
        return error;
    }

    return test_cl_image_read(context, queue, target, image, width, height,
                              depth, sampleNum, outFormat, outType,
                              outResultBuffer);
}

static int test_image_format_read(cl_context context, cl_command_queue queue,
                                  size_t width, size_t height, size_t depth,
                                  GLenum target, const format *fmt, MTdata data)
{
    int error = 0;

    // Determine the maximum number of supported samples
    GLint samples = 1;
    if (target == GL_TEXTURE_2D_MULTISAMPLE
        || target == GL_TEXTURE_2D_MULTISAMPLE_ARRAY)
        samples = get_gl_max_samples(target, fmt->internal);

    // If we're testing a half float format, then we need to determine the
    // rounding mode of this machine.  Punt if we fail to do so.

    if (fmt->type == kHalf)
    {
        if (DetectFloatToHalfRoundingMode(queue)) return 1;
        bool supports_half = false;
        error = supportsHalf(context, &supports_half);
        if (error != 0) return error;
        if (!supports_half) return 0;
    }
#ifdef GL_VERSION_3_2
    if (get_base_gl_target(target) == GL_TEXTURE_2D_MULTISAMPLE
        || get_base_gl_target(target) == GL_TEXTURE_2D_MULTISAMPLE_ARRAY)
    {
        bool supports_msaa;
        error = supportsMsaa(context, &supports_msaa);
        if (error != 0) return error;
        if (!supports_msaa) return 0;
    }
    if (fmt->formattype == GL_DEPTH_COMPONENT
        || fmt->formattype == GL_DEPTH_STENCIL)
    {
        bool supports_depth;
        error = supportsDepth(context, &supports_depth);
        if (error != 0) return error;
        if (!supports_depth) return 0;
    }
#endif
    size_t w = width, h = height, d = depth;

    // Unpack the format and use it, along with the target, to create an
    // appropriate GL texture.

    GLenum gl_fmt = fmt->formattype;
    GLenum gl_internal_fmt = fmt->internal;
    GLenum gl_type = fmt->datatype;
    ExplicitType type = fmt->type;

    // Required for most of the texture-backed cases:
    glTextureWrapper texture;

    // Required for the special case of TextureBuffer textures:
    glBufferWrapper glbuf;

    // And these are required for the case of Renderbuffer images:
    glFramebufferWrapper glFramebuffer;
    glRenderbufferWrapper glRenderbuffer;

    void *buffer = NULL;

    // Use the correct texture creation function depending on the target, and
    // adjust width, height, depth as appropriate so subsequent size
    // calculations succeed.

    switch (get_base_gl_target(target))
    {
        case GL_TEXTURE_1D:
            h = 1;
            d = 1;
            buffer =
                CreateGLTexture1D(width, target, gl_fmt, gl_internal_fmt,
                                  gl_type, type, &texture, &error, true, data);
            break;
        case GL_TEXTURE_BUFFER:
            h = 1;
            d = 1;
            buffer = CreateGLTextureBuffer(
                width, target, gl_fmt, gl_internal_fmt, gl_type, type, &texture,
                &glbuf, &error, true, data);
            break;
        case GL_RENDERBUFFER:
        case GL_COLOR_ATTACHMENT0:
            d = 1;
            buffer = CreateGLRenderbuffer(
                width, height, target, gl_fmt, gl_internal_fmt, gl_type, type,
                &glFramebuffer, &glRenderbuffer, &error, data, true);
            break;
        case GL_TEXTURE_2D:
        case GL_TEXTURE_RECTANGLE_EXT:
        case GL_TEXTURE_CUBE_MAP:
            d = 1;
            buffer = CreateGLTexture2D(width, height, target, gl_fmt,
                                       gl_internal_fmt, gl_type, type, &texture,
                                       &error, true, data);
            break;
        case GL_TEXTURE_1D_ARRAY:
            d = 1;
            buffer = CreateGLTexture1DArray(width, height, target, gl_fmt,
                                            gl_internal_fmt, gl_type, type,
                                            &texture, &error, true, data);
            break;
        case GL_TEXTURE_2D_ARRAY:
            buffer = CreateGLTexture2DArray(width, height, depth, target,
                                            gl_fmt, gl_internal_fmt, gl_type,
                                            type, &texture, &error, true, data);
            break;
        case GL_TEXTURE_3D:
            buffer = CreateGLTexture3D(width, height, depth, target, gl_fmt,
                                       gl_internal_fmt, gl_type, type, &texture,
                                       &error, data, true);
            break;
#ifdef GL_VERSION_3_2
        case GL_TEXTURE_2D_MULTISAMPLE:
            d = 1;
            buffer = CreateGLTexture2DMultisample(
                width, height, samples, target, gl_fmt, gl_internal_fmt,
                gl_type, type, &texture, &error, true, data, true);
            break;
        case GL_TEXTURE_2D_MULTISAMPLE_ARRAY:
            buffer = CreateGLTexture2DArrayMultisample(
                width, height, depth, samples, target, gl_fmt, gl_internal_fmt,
                gl_type, type, &texture, &error, true, data, true);
            break;
#endif
        default: log_error("Unsupported texture target."); return 1;
    }

    if (error == -2)
    {
        log_info("OpenGL texture couldn't be created, because a texture is too "
                 "big. Skipping test.\n");
        return 0;
    }

    // Check to see if the texture could not be created for some other reason
    // like GL_FRAMEBUFFER_UNSUPPORTED
    if (error == GL_FRAMEBUFFER_UNSUPPORTED)
    {
        log_info("Skipping...\n");
        return 0;
    }

    if (error != 0)
    {
        if ((gl_fmt == GL_RGBA_INTEGER_EXT)
            && (!CheckGLIntegerExtensionSupport()))
        {
            log_info("OpenGL version does not support GL_RGBA_INTEGER_EXT. "
                     "Skipping test.\n");
            return 0;
        }
        else
        {
            return error;
        }
    }

    BufferOwningPtr<char> inputBuffer(buffer);
    if (inputBuffer == NULL) return -1;

    cl_image_format clFormat;
    ExplicitType actualType;
    char *outBuffer;

    // Perform the read:

    GLuint globj = texture;
    if (target == GL_RENDERBUFFER || target == GL_COLOR_ATTACHMENT0)
    {
        globj = glRenderbuffer;
    }

    error = test_image_read(context, queue, target, globj, w, h, d, samples,
                            &clFormat, &actualType, (void **)&outBuffer);

    if (error != 0) return error;

    BufferOwningPtr<char> actualResults(outBuffer);
    if (actualResults == NULL) return -1;

    log_info("- Read [%4d x %4d x %4d x %4d] : GL Texture : %s : %s : %s => CL "
             "Image : %s : %s \n",
             (int)w, (int)h, (int)d, (int)samples, GetGLFormatName(gl_fmt),
             GetGLFormatName(gl_internal_fmt), GetGLTypeName(gl_type),
             GetChannelOrderName(clFormat.image_channel_order),
             GetChannelTypeName(clFormat.image_channel_data_type));

    BufferOwningPtr<char> convertedInputs;

    // We have to convert our input buffer to the returned type, so we can
    // validate. This is necessary because OpenCL might not actually pick an
    // internal format that actually matches our input format (for example, if
    // it picks a normalized format, the results will come out as floats instead
    // of going in as ints).

    if (gl_type == GL_UNSIGNED_INT_2_10_10_10_REV)
    {
        cl_uint *p = (cl_uint *)buffer;
        float *inData = (float *)malloc(w * h * d * samples * sizeof(float));

        for (size_t i = 0; i < 4 * w * h * d * samples; i += 4)
        {
            inData[i + 0] = (float)((p[0] >> 20) & 0x3ff) / (float)1023;
            inData[i + 1] = (float)((p[0] >> 10) & 0x3ff) / (float)1023;
            inData[i + 2] = (float)(p[0] & 0x3ff) / (float)1023;
            p++;
        }

        convertedInputs.reset(inData);
        if (convertedInputs == NULL) return -1;
    }
    else if (gl_type == GL_DEPTH24_STENCIL8)
    {
        // GL_DEPTH24_STENCIL8 is treated as CL_UNORM_INT24 + CL_DEPTH_STENCIL
        // where the stencil is ignored.
        cl_uint *p = (cl_uint *)buffer;
        float *inData = (float *)malloc(w * h * d * samples * sizeof(float));

        for (size_t i = 0; i < w * h * d * samples; i++)
        {
            inData[i] = (float)((p[i] >> 8) & 0xffffff) / (float)0xfffffe;
        }

        convertedInputs.reset(inData);
        if (convertedInputs == NULL) return -1;
    }
    else if (gl_type == GL_FLOAT_32_UNSIGNED_INT_24_8_REV)
    {
        // GL_FLOAT_32_UNSIGNED_INT_24_8_REV is treated as a CL_FLOAT +
        // unused 24 + CL_DEPTH_STENCIL; we check the float value and ignore the
        // second word

        float *p = (float *)buffer;
        float *inData = (float *)malloc(w * h * d * samples * sizeof(float));

        for (size_t i = 0; i < w * h * d * samples; i++)
        {
            inData[i] = p[i * 2];
        }

        convertedInputs.reset(inData);
        if (convertedInputs == NULL) return -1;
    }
    else
    {
        convertedInputs.reset(convert_to_expected(
            inputBuffer, w * h * d * samples, type, actualType,
            get_channel_order_channel_count(clFormat.image_channel_order)));
        if (convertedInputs == NULL) return -1;
    }

    // Now we validate
    if (actualType == kFloat)
    {
        if (clFormat.image_channel_data_type == CL_UNORM_INT_101010)
        {
            return validate_float_results_rgb_101010(
                convertedInputs, actualResults, w, h, d, samples);
        }
        else
        {
            return validate_float_results(
                convertedInputs, actualResults, w, h, d, samples,
                get_channel_order_channel_count(clFormat.image_channel_order));
        }
    }
    else
    {
        return validate_integer_results(convertedInputs, actualResults, w, h, d,
                                        samples,
                                        get_explicit_type_size(actualType));
    }
}

int test_images_read_common(cl_device_id device, cl_context context,
                            cl_command_queue queue, const format *formats,
                            size_t nformats, GLenum *targets, size_t ntargets,
                            sizevec_t *sizes, size_t nsizes)
{
    int error = 0;
    RandomSeed seed(gRandomSeed);

    // First, ensure this device supports images.

    if (checkForImageSupport(device))
    {
        log_info("Device does not support images.  Skipping test.\n");
        return 0;
    }

    size_t fidx, tidx, sidx;

    // Test each format on every target, every size.

    for (fidx = 0; fidx < nformats; fidx++)
    {
        for (tidx = 0; tidx < ntargets; tidx++)
        {

            // Texture buffer only takes an internal format, so the level data
            // passed by the test and used for verification must match the
            // internal format
            if ((targets[tidx] == GL_TEXTURE_BUFFER)
                && (GetGLFormat(formats[fidx].internal)
                    != formats[fidx].formattype))
                continue;

            if (formats[fidx].datatype == GL_UNSIGNED_INT_2_10_10_10_REV)
            {
                // Check if the RGB 101010 format is supported
                if (is_rgb_101010_supported(context, targets[tidx]) == 0)
                    break; // skip
            }

            if (targets[tidx] != GL_TEXTURE_BUFFER)
                log_info("Testing image read for GL format %s : %s : %s : %s\n",
                         GetGLTargetName(targets[tidx]),
                         GetGLFormatName(formats[fidx].internal),
                         GetGLBaseFormatName(formats[fidx].formattype),
                         GetGLTypeName(formats[fidx].datatype));
            else
                log_info("Testing image read for GL format %s : %s\n",
                         GetGLTargetName(targets[tidx]),
                         GetGLFormatName(formats[fidx].internal));

            for (sidx = 0; sidx < nsizes; sidx++)
            {

                // Test this format + size:
                int err;
                if ((err = test_image_format_read(
                         context, queue, sizes[sidx].width, sizes[sidx].height,
                         sizes[sidx].depth, targets[tidx], &formats[fidx],
                         seed)))
                {
                    // Negative return values are errors, positive mean the test
                    // was skipped
                    if (err < 0)
                    {

                        // We land here in the event of test failure.

                        log_error("ERROR: Image read test failed for %s : %s : "
                                  "%s : %s\n\n",
                                  GetGLTargetName(targets[tidx]),
                                  GetGLFormatName(formats[fidx].internal),
                                  GetGLBaseFormatName(formats[fidx].formattype),
                                  GetGLTypeName(formats[fidx].datatype));
                        error++;
                    }

                    // Skip the other sizes for this format.
                    printf("Skipping remaining sizes for this format\n");

                    break;
                }
            }

            // Note a successful format test, if we passed every size.

            if (sidx == nsizes)
            {
                log_info("passed: Image read test for GL format  %s : %s : %s "
                         ": %s\n\n",
                         GetGLTargetName(targets[tidx]),
                         GetGLFormatName(formats[fidx].internal),
                         GetGLBaseFormatName(formats[fidx].formattype),
                         GetGLTypeName(formats[fidx].datatype));
            }
        }
    }

    return error;
}