| // |
| // 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); |
| |
| 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"; |
| |
| 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, struct 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, struct 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 == sizeof (sizes) / sizeof( sizes[0] ) ) { |
| 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; |
| } |