// // 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 #else #include #include #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 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 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 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; }