mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-19 06:09:01 +00:00
53db6e7f9f
The maintenance of the conformance tests is moving to Github. This commit contains all the changes that have been done in Gitlab since the first public release of the conformance tests. Signed-off-by: Kevin Petit <kevin.petit@arm.com>
869 lines
30 KiB
C++
869 lines
30 KiB
C++
//
|
|
// 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 "testBase.h"
|
|
#include "common.h"
|
|
#include <limits.h>
|
|
|
|
#if defined( __APPLE__ )
|
|
#include <OpenGL/glu.h>
|
|
#else
|
|
#include <GL/glu.h>
|
|
#include <CL/cl_gl.h>
|
|
#endif
|
|
|
|
extern "C" { extern cl_uint gRandomSeed; };
|
|
|
|
#pragma mark -
|
|
#pragma mark Write test kernels
|
|
|
|
static const char *kernelpattern_image_write_1D =
|
|
"__kernel void sample_test( __global %s4 *source, write_only image1d_t dest )\n"
|
|
"{\n"
|
|
" uint index = get_global_id(0);\n"
|
|
" %s4 value = source[index];\n"
|
|
" write_image%s( dest, index, %s(value));\n"
|
|
"}\n";
|
|
|
|
static const char *kernelpattern_image_write_1D_half =
|
|
"__kernel void sample_test( __global half4 *source, write_only image1d_t dest )\n"
|
|
"{\n"
|
|
" uint index = get_global_id(0);\n"
|
|
" write_imagef( dest, index, vload_half4(index, (__global half *)source));\n"
|
|
"}\n";
|
|
|
|
static const char *kernelpattern_image_write_1D_buffer =
|
|
"__kernel void sample_test( __global %s4 *source, write_only image1d_buffer_t dest )\n"
|
|
"{\n"
|
|
" uint index = get_global_id(0);\n"
|
|
" %s4 value = source[index];\n"
|
|
" write_image%s( dest, index, %s(value));\n"
|
|
"}\n";
|
|
|
|
static const char *kernelpattern_image_write_1D_buffer_half =
|
|
"__kernel void sample_test( __global half4 *source, write_only image1d_buffer_t dest )\n"
|
|
"{\n"
|
|
" uint index = get_global_id(0);\n"
|
|
" write_imagef( dest, index, vload_half4(index, (__global half *)source));\n"
|
|
"}\n";
|
|
|
|
static const char *kernelpattern_image_write_2D =
|
|
"__kernel void sample_test( __global %s4 *source, write_only image2d_t dest )\n"
|
|
"{\n"
|
|
" int tidX = get_global_id(0);\n"
|
|
" int tidY = get_global_id(1);\n"
|
|
" uint index = tidY * get_image_width( dest ) + tidX;\n"
|
|
" %s4 value = source[index];\n"
|
|
" write_image%s( dest, (int2)( tidX, tidY ), %s(value));\n"
|
|
"}\n";
|
|
|
|
static const char *kernelpattern_image_write_2D_half =
|
|
"__kernel void sample_test( __global half4 *source, write_only image2d_t dest )\n"
|
|
"{\n"
|
|
" int tidX = get_global_id(0);\n"
|
|
" int tidY = get_global_id(1);\n"
|
|
" uint index = tidY * get_image_width( dest ) + tidX;\n"
|
|
" write_imagef( dest, (int2)( tidX, tidY ), vload_half4(index, (__global half *)source));\n"
|
|
"}\n";
|
|
|
|
static const char *kernelpattern_image_write_1Darray =
|
|
"__kernel void sample_test( __global %s4 *source, write_only image1d_array_t dest )\n"
|
|
"{\n"
|
|
" int tidX = get_global_id(0);\n"
|
|
" int tidY = get_global_id(1);\n"
|
|
" uint index = tidY * get_image_width( dest ) + tidX;\n"
|
|
" %s4 value = source[index];\n"
|
|
" write_image%s( dest, (int2)( tidX, tidY ), %s(value));\n"
|
|
"}\n";
|
|
|
|
static const char *kernelpattern_image_write_1Darray_half =
|
|
"__kernel void sample_test( __global half4 *source, write_only image1d_array_t dest )\n"
|
|
"{\n"
|
|
" int tidX = get_global_id(0);\n"
|
|
" int tidY = get_global_id(1);\n"
|
|
" uint index = tidY * get_image_width( dest ) + tidX;\n"
|
|
" write_imagef( dest, (int2)( tidX, tidY ), vload_half4(index, (__global half *)source));\n"
|
|
"}\n";
|
|
|
|
static const char *kernelpattern_image_write_3D =
|
|
"__kernel void sample_test( __global %s4 *source, write_only image3d_t dest )\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( dest );\n"
|
|
" int height = get_image_height( dest );\n"
|
|
" int index = tidZ * width * height + tidY * width + tidX;\n"
|
|
" %s4 value = source[index];\n"
|
|
" write_image%s( dest, (int4)( tidX, tidY, tidZ, 0 ), %s(value));\n"
|
|
"}\n";
|
|
|
|
static const char *kernelpattern_image_write_3D_half =
|
|
"__kernel void sample_test( __global half4 *source, write_only image3d_t dest )\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( dest );\n"
|
|
" int height = get_image_height( dest );\n"
|
|
" int index = tidZ * width * height + tidY * width + tidX;\n"
|
|
" write_imagef( dest, (int4)( tidX, tidY, tidZ, 0 ), vload_half4(index, (__global half *)source));\n"
|
|
"}\n";
|
|
|
|
static const char *kernelpattern_image_write_2Darray =
|
|
"__kernel void sample_test( __global %s4 *source, write_only image2d_array_t dest )\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( dest );\n"
|
|
" int height = get_image_height( dest );\n"
|
|
" int index = tidZ * width * height + tidY * width + tidX;\n"
|
|
" %s4 value = source[index];\n"
|
|
" write_image%s( dest, (int4)( tidX, tidY, tidZ, 0 ), %s(value));\n"
|
|
"}\n";
|
|
|
|
static const char *kernelpattern_image_write_2Darray_half =
|
|
"__kernel void sample_test( __global half4 *source, write_only image2d_array_t dest )\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( dest );\n"
|
|
" int height = get_image_height( dest );\n"
|
|
" int index = tidZ * width * height + tidY * width + tidX;\n"
|
|
" write_imagef( dest, (int4)( tidX, tidY, tidZ, 0 ), vload_half4(index, (__global half *)source));\n"
|
|
"}\n";
|
|
|
|
#ifdef GL_VERSION_3_2
|
|
|
|
static const char * kernelpattern_image_write_2D_depth =
|
|
"__kernel void sample_test( __global %s *source, write_only image2d_depth_t dest )\n"
|
|
"{\n"
|
|
" int tidX = get_global_id(0);\n"
|
|
" int tidY = get_global_id(1);\n"
|
|
" uint index = tidY * get_image_width( dest ) + tidX;\n"
|
|
" float value = source[index];\n"
|
|
" write_imagef( dest, (int2)( tidX, tidY ), value);\n"
|
|
"}\n";
|
|
|
|
static const char * kernelpattern_image_write_2D_array_depth =
|
|
"__kernel void sample_test( __global %s *source, write_only image2d_array_depth_t dest )\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( dest );\n"
|
|
" int height = get_image_height( dest );\n"
|
|
" int index = tidZ * width * height + tidY * width + tidX;\n"
|
|
" %s value = source[index];\n"
|
|
" write_image%s( dest, (int4)( tidX, tidY, tidZ, 0 ), %s(value));\n"
|
|
"}\n";
|
|
|
|
|
|
#endif
|
|
|
|
#pragma mark -
|
|
#pragma mark Utility functions
|
|
|
|
static const char* get_appropriate_write_kernel(GLenum target,
|
|
ExplicitType type, cl_channel_order channel_order)
|
|
{
|
|
switch (get_base_gl_target(target)) {
|
|
case GL_TEXTURE_1D:
|
|
|
|
if (type == kHalf)
|
|
return kernelpattern_image_write_1D_half;
|
|
else
|
|
return kernelpattern_image_write_1D;
|
|
break;
|
|
case GL_TEXTURE_BUFFER:
|
|
if (type == kHalf)
|
|
return kernelpattern_image_write_1D_buffer_half;
|
|
else
|
|
return kernelpattern_image_write_1D_buffer;
|
|
break;
|
|
case GL_TEXTURE_1D_ARRAY:
|
|
if (type == kHalf)
|
|
return kernelpattern_image_write_1Darray_half;
|
|
else
|
|
return kernelpattern_image_write_1Darray;
|
|
break;
|
|
case GL_COLOR_ATTACHMENT0:
|
|
case GL_RENDERBUFFER:
|
|
case GL_TEXTURE_RECTANGLE_EXT:
|
|
case GL_TEXTURE_2D:
|
|
case GL_TEXTURE_CUBE_MAP:
|
|
#ifdef GL_VERSION_3_2
|
|
if (channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
|
|
return kernelpattern_image_write_2D_depth;
|
|
#endif
|
|
if (type == kHalf)
|
|
return kernelpattern_image_write_2D_half;
|
|
else
|
|
return kernelpattern_image_write_2D;
|
|
break;
|
|
|
|
case GL_TEXTURE_2D_ARRAY:
|
|
#ifdef GL_VERSION_3_2
|
|
if (channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
|
|
return kernelpattern_image_write_2D_array_depth;
|
|
#endif
|
|
if (type == kHalf)
|
|
return kernelpattern_image_write_2Darray_half;
|
|
else
|
|
return kernelpattern_image_write_2Darray;
|
|
break;
|
|
|
|
case GL_TEXTURE_3D:
|
|
if (type == kHalf)
|
|
return kernelpattern_image_write_3D_half;
|
|
else
|
|
return kernelpattern_image_write_3D;
|
|
break;
|
|
|
|
default:
|
|
log_error("Unsupported GL tex target (%s) passed to write test: "
|
|
"%s (%s):%d", GetGLTargetName(target), __FUNCTION__,
|
|
__FILE__, __LINE__);
|
|
return NULL;
|
|
}
|
|
}
|
|
|
|
void set_dimensions_by_target(GLenum target, size_t *dims, size_t sizes[3],
|
|
size_t width, size_t height, size_t depth)
|
|
{
|
|
switch (get_base_gl_target(target)) {
|
|
case GL_TEXTURE_1D:
|
|
sizes[0] = width;
|
|
*dims = 1;
|
|
break;
|
|
|
|
case GL_TEXTURE_BUFFER:
|
|
sizes[0] = width;
|
|
*dims = 1;
|
|
break;
|
|
|
|
case GL_TEXTURE_1D_ARRAY:
|
|
sizes[0] = width;
|
|
sizes[1] = height;
|
|
*dims = 2;
|
|
break;
|
|
|
|
case GL_COLOR_ATTACHMENT0:
|
|
case GL_RENDERBUFFER:
|
|
case GL_TEXTURE_RECTANGLE_EXT:
|
|
case GL_TEXTURE_2D:
|
|
case GL_TEXTURE_CUBE_MAP:
|
|
|
|
sizes[0] = width;
|
|
sizes[1] = height;
|
|
*dims = 2;
|
|
break;
|
|
|
|
case GL_TEXTURE_2D_ARRAY:
|
|
sizes[0] = width;
|
|
sizes[1] = height;
|
|
sizes[2] = depth;
|
|
*dims = 3;
|
|
break;
|
|
|
|
case GL_TEXTURE_3D:
|
|
sizes[0] = width;
|
|
sizes[1] = height;
|
|
sizes[2] = depth;
|
|
*dims = 3;
|
|
break;
|
|
|
|
default:
|
|
log_error("Unsupported GL tex target (%s) passed to write test: "
|
|
"%s (%s):%d", GetGLTargetName(target), __FUNCTION__,
|
|
__FILE__, __LINE__);
|
|
}
|
|
}
|
|
|
|
int test_cl_image_write( cl_context context, cl_command_queue queue,
|
|
GLenum target, cl_mem clImage, size_t width, size_t height, size_t depth,
|
|
cl_image_format *outFormat, ExplicitType *outType, void **outSourceBuffer,
|
|
MTdata d, bool supports_half )
|
|
{
|
|
size_t global_dims, global_sizes[3];
|
|
clProgramWrapper program;
|
|
clKernelWrapper kernel;
|
|
clMemWrapper inStream;
|
|
char* programPtr;
|
|
int error;
|
|
char kernelSource[2048];
|
|
|
|
// What CL format did we get from the texture?
|
|
|
|
error = clGetImageInfo(clImage, CL_IMAGE_FORMAT, sizeof(cl_image_format),
|
|
outFormat, NULL);
|
|
test_error(error, "Unable to get the CL image format");
|
|
|
|
// Create the kernel source. The target and the data type will influence
|
|
// which particular kernel we choose.
|
|
|
|
*outType = get_write_kernel_type( outFormat );
|
|
size_t channelSize = get_explicit_type_size(*outType);
|
|
|
|
const char* appropriateKernel = get_appropriate_write_kernel(target,
|
|
*outType, outFormat->image_channel_order);
|
|
if (*outType == kHalf && !supports_half) {
|
|
log_info("cl_khr_fp16 isn't supported. Skip this test.\n");
|
|
return 0;
|
|
}
|
|
|
|
const char* suffix = get_kernel_suffix( outFormat );
|
|
const char* convert = get_write_conversion( outFormat, *outType );
|
|
|
|
sprintf(kernelSource, appropriateKernel, get_explicit_type_name( *outType ),
|
|
get_explicit_type_name( *outType ), suffix, convert);
|
|
|
|
programPtr = kernelSource;
|
|
if( create_single_kernel_helper_with_build_options( context, &program, &kernel, 1,
|
|
(const char **)&programPtr, "sample_test", "-cl-std=CL2.0" ) )
|
|
{
|
|
return -1;
|
|
}
|
|
|
|
// Create an appropriately-sized output buffer.
|
|
|
|
// Check to see if the output buffer will fit on the device
|
|
size_t bytes = channelSize * 4 * width * height * depth;
|
|
cl_ulong alloc_size = 0;
|
|
|
|
cl_device_id device = NULL;
|
|
error = clGetCommandQueueInfo(queue, CL_QUEUE_DEVICE, sizeof(device), &device, NULL);
|
|
test_error( error, "Unable to query command queue for device" );
|
|
|
|
error = clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(alloc_size), &alloc_size, NULL);
|
|
test_error( error, "Unable to device for max mem alloc size" );
|
|
|
|
if (bytes > alloc_size) {
|
|
log_info(" Skipping: Buffer size (%lu) is greater than CL_DEVICE_MAX_MEM_ALLOC_SIZE (%lu)\n", bytes, alloc_size);
|
|
*outSourceBuffer = NULL;
|
|
return 0;
|
|
}
|
|
|
|
*outSourceBuffer = CreateRandomData(*outType, width * height * depth * 4, d);
|
|
|
|
inStream = clCreateBuffer( context, CL_MEM_COPY_HOST_PTR,
|
|
channelSize * 4 * width * height * depth, *outSourceBuffer, &error );
|
|
test_error( error, "Unable to create output buffer" );
|
|
|
|
cl_sampler_properties properties[] = {
|
|
CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE,
|
|
CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_NONE,
|
|
CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST,
|
|
0 };
|
|
clSamplerWrapper sampler = clCreateSamplerWithProperties( context, properties, &error );
|
|
test_error( error, "Unable to create sampler" );
|
|
|
|
error = clSetKernelArg( kernel, 0, sizeof( inStream ), &inStream );
|
|
test_error( error, "Unable to set kernel arguments" );
|
|
|
|
error = clSetKernelArg( kernel, 1, sizeof( clImage ), &clImage );
|
|
test_error( error, "Unable to set kernel arguments" );
|
|
|
|
// Flush and Acquire.
|
|
|
|
glFinish();
|
|
|
|
error = (*clEnqueueAcquireGLObjects_ptr)( queue, 1, &clImage, 0, NULL, NULL);
|
|
test_error( error, "Unable to acquire GL obejcts");
|
|
|
|
// Execute ( letting OpenCL choose the local size )
|
|
|
|
// Setup the global dimensions and sizes based on the target type.
|
|
set_dimensions_by_target(target, &global_dims, global_sizes,
|
|
width, height, depth);
|
|
|
|
error = clEnqueueNDRangeKernel( queue, kernel, global_dims, NULL,
|
|
global_sizes, NULL, 0, NULL, NULL );
|
|
test_error( error, "Unable to execute test kernel" );
|
|
|
|
clEventWrapper event;
|
|
error = (*clEnqueueReleaseGLObjects_ptr)( queue, 1, &clImage, 0, NULL, &event );
|
|
test_error(error, "clEnqueueReleaseGLObjects failed");
|
|
|
|
error = clWaitForEvents( 1, &event );
|
|
test_error(error, "clWaitForEvents failed");
|
|
|
|
return 0;
|
|
}
|
|
|
|
static int test_image_write( cl_context context, cl_command_queue queue,
|
|
GLenum glTarget, GLuint glTexture, size_t width, size_t height, size_t depth,
|
|
cl_image_format *outFormat, ExplicitType *outType, void **outSourceBuffer,
|
|
MTdata d, bool supports_half )
|
|
{
|
|
int error;
|
|
|
|
// Create a CL image from the supplied GL texture
|
|
clMemWrapper image = (*clCreateFromGLTexture_ptr)( context, CL_MEM_WRITE_ONLY,
|
|
glTarget, 0, glTexture, &error );
|
|
|
|
if ( error != CL_SUCCESS ) {
|
|
print_error( error, "Unable to create CL image from GL texture" );
|
|
GLint fmt;
|
|
glGetTexLevelParameteriv( glTarget, 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_write( context, queue, glTarget, image,
|
|
width, height, depth, outFormat, outType, outSourceBuffer, d, supports_half );
|
|
}
|
|
|
|
int supportsHalf(cl_context context, bool* supports_half)
|
|
{
|
|
int error;
|
|
size_t size;
|
|
cl_uint numDev;
|
|
|
|
error = clGetContextInfo(context, CL_CONTEXT_NUM_DEVICES, sizeof(cl_uint), &numDev, NULL);
|
|
test_error(error, "clGetContextInfo for CL_CONTEXT_NUM_DEVICES failed");
|
|
|
|
cl_device_id* devices = new cl_device_id[numDev];
|
|
error = clGetContextInfo(context, CL_CONTEXT_DEVICES, numDev * sizeof(cl_device_id), devices, NULL);
|
|
test_error(error, "clGetContextInfo for CL_CONTEXT_DEVICES failed");
|
|
|
|
// Get the extensions string for the device
|
|
error = clGetDeviceInfo(devices[0], CL_DEVICE_EXTENSIONS, 0, NULL, &size);
|
|
test_error(error, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS size failed");
|
|
|
|
char *extensions = new char[size+1];
|
|
if (extensions == 0) {
|
|
log_error("Failed to allocate memory for extensions string.\n");
|
|
return -1;
|
|
}
|
|
memset( extensions, CHAR_MIN, sizeof(char)*(size+1) );
|
|
|
|
error = clGetDeviceInfo(devices[0], CL_DEVICE_EXTENSIONS, sizeof(char)*size, extensions, NULL);
|
|
test_error(error, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS failed");
|
|
|
|
*supports_half = strstr(extensions, "cl_khr_fp16");
|
|
delete [] extensions;
|
|
delete [] devices;
|
|
|
|
return error;
|
|
}
|
|
|
|
int supportsMsaa(cl_context context, bool* supports_msaa)
|
|
{
|
|
int error;
|
|
size_t size;
|
|
cl_uint numDev;
|
|
|
|
error = clGetContextInfo(context, CL_CONTEXT_NUM_DEVICES, sizeof(cl_uint), &numDev, NULL);
|
|
test_error(error, "clGetContextInfo for CL_CONTEXT_NUM_DEVICES failed");
|
|
|
|
cl_device_id* devices = new cl_device_id[numDev];
|
|
error = clGetContextInfo(context, CL_CONTEXT_DEVICES, numDev * sizeof(cl_device_id), devices, NULL);
|
|
test_error(error, "clGetContextInfo for CL_CONTEXT_DEVICES failed");
|
|
|
|
// Get the extensions string for the device
|
|
error = clGetDeviceInfo(devices[0], CL_DEVICE_EXTENSIONS, 0, NULL, &size);
|
|
test_error(error, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS size failed");
|
|
|
|
char *extensions = new char[size+1];
|
|
if (extensions == 0) {
|
|
log_error("Failed to allocate memory for extensions string.\n");
|
|
return -1;
|
|
}
|
|
memset( extensions, CHAR_MIN, sizeof(char)*(size+1) );
|
|
|
|
error = clGetDeviceInfo(devices[0], CL_DEVICE_EXTENSIONS, sizeof(char)*size, extensions, NULL);
|
|
test_error(error, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS failed");
|
|
|
|
*supports_msaa = strstr(extensions, "cl_khr_gl_msaa_sharing");
|
|
delete [] extensions;
|
|
delete [] devices;
|
|
|
|
return error;
|
|
}
|
|
|
|
int supportsDepth(cl_context context, bool* supports_depth)
|
|
{
|
|
int error;
|
|
size_t size;
|
|
cl_uint numDev;
|
|
|
|
error = clGetContextInfo(context, CL_CONTEXT_NUM_DEVICES, sizeof(cl_uint), &numDev, NULL);
|
|
test_error(error, "clGetContextInfo for CL_CONTEXT_NUM_DEVICES failed");
|
|
|
|
cl_device_id* devices = new cl_device_id[numDev];
|
|
error = clGetContextInfo(context, CL_CONTEXT_DEVICES, numDev * sizeof(cl_device_id), devices, NULL);
|
|
test_error(error, "clGetContextInfo for CL_CONTEXT_DEVICES failed");
|
|
|
|
// Get the extensions string for the device
|
|
error = clGetDeviceInfo(devices[0], CL_DEVICE_EXTENSIONS, 0, NULL, &size);
|
|
test_error(error, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS size failed");
|
|
|
|
char *extensions = new char[size+1];
|
|
if (extensions == 0) {
|
|
log_error("Failed to allocate memory for extensions string.\n");
|
|
return -1;
|
|
}
|
|
memset( extensions, CHAR_MIN, sizeof(char)*(size+1) );
|
|
|
|
error = clGetDeviceInfo(devices[0], CL_DEVICE_EXTENSIONS, sizeof(char)*size, extensions, NULL);
|
|
test_error(error, "clGetDeviceInfo for CL_DEVICE_EXTENSIONS failed");
|
|
|
|
*supports_depth = strstr(extensions, "cl_khr_gl_depth_images");
|
|
delete [] extensions;
|
|
delete [] devices;
|
|
|
|
return error;
|
|
}
|
|
|
|
static int test_image_format_write( cl_context context, cl_command_queue queue,
|
|
size_t width, size_t height, size_t depth, GLenum target, GLenum format,
|
|
GLenum internalFormat, GLenum glType, ExplicitType type, MTdata d )
|
|
{
|
|
int error;
|
|
int samples = 8;
|
|
// 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( type == kHalf )
|
|
if( DetectFloatToHalfRoundingMode(queue) )
|
|
return 1;
|
|
|
|
// Create an appropriate GL texture or renderbuffer, given the target.
|
|
|
|
glTextureWrapper glTexture;
|
|
glBufferWrapper glBuf;
|
|
glFramebufferWrapper glFramebuffer;
|
|
glRenderbufferWrapper glRenderbuffer;
|
|
switch (get_base_gl_target(target)) {
|
|
case GL_TEXTURE_1D:
|
|
CreateGLTexture1D( width, target, format, internalFormat, glType,
|
|
type, &glTexture, &error, false, d );
|
|
break;
|
|
case GL_TEXTURE_BUFFER:
|
|
CreateGLTextureBuffer( width, target, format, internalFormat, glType,
|
|
type, &glTexture, &glBuf, &error, false, d );
|
|
break;
|
|
case GL_TEXTURE_1D_ARRAY:
|
|
CreateGLTexture1DArray( width, height, target, format, internalFormat,
|
|
glType, type, &glTexture, &error, false, d );
|
|
break;
|
|
case GL_TEXTURE_RECTANGLE_EXT:
|
|
case GL_TEXTURE_2D:
|
|
case GL_TEXTURE_CUBE_MAP:
|
|
CreateGLTexture2D( width, height, target, format, internalFormat, glType,
|
|
type, &glTexture, &error, false, d );
|
|
break;
|
|
case GL_COLOR_ATTACHMENT0:
|
|
case GL_RENDERBUFFER:
|
|
CreateGLRenderbuffer(width, height, target, format, internalFormat,
|
|
glType, type, &glFramebuffer, &glRenderbuffer, &error, d, false);
|
|
case GL_TEXTURE_2D_ARRAY:
|
|
CreateGLTexture2DArray( width, height, depth, target, format,
|
|
internalFormat, glType, type, &glTexture, &error, false, d );
|
|
break;
|
|
case GL_TEXTURE_3D:
|
|
CreateGLTexture3D( width, height, depth, target, format,
|
|
internalFormat, glType, type, &glTexture, &error, d, false );
|
|
break;
|
|
|
|
default:
|
|
log_error("Unsupported GL tex target (%s) passed to write test: "
|
|
"%s (%s):%d", GetGLTargetName(target), __FUNCTION__,
|
|
__FILE__, __LINE__);
|
|
}
|
|
|
|
// If there was a problem during creation, make sure it isn't a known
|
|
// cause, and then complain.
|
|
if ( error == -2 ) {
|
|
log_info("OpenGL texture couldn't be created, because a texture is too big. Skipping test.\n");
|
|
return 0;
|
|
}
|
|
|
|
if ( error != 0 ) {
|
|
if ((format == GL_RGBA_INTEGER_EXT) && (!CheckGLIntegerExtensionSupport())){
|
|
log_info("OpenGL version does not support GL_RGBA_INTEGER_EXT. "
|
|
"Skipping test.\n");
|
|
return 0;
|
|
} else {
|
|
return error;
|
|
}
|
|
}
|
|
|
|
// Run and get the results
|
|
cl_image_format clFormat;
|
|
ExplicitType sourceType;
|
|
ExplicitType validationType;
|
|
void *outSourceBuffer = NULL;
|
|
|
|
GLenum globj = glTexture;
|
|
if (target == GL_RENDERBUFFER || target == GL_COLOR_ATTACHMENT0) {
|
|
globj = glRenderbuffer;
|
|
}
|
|
|
|
bool supports_half = false;
|
|
error = supportsHalf(context, &supports_half);
|
|
if( error != 0 )
|
|
return error;
|
|
|
|
error = test_image_write( context, queue, target, globj, width, height,
|
|
depth, &clFormat, &sourceType, (void **)&outSourceBuffer, d, supports_half );
|
|
|
|
if( error != 0 || ((sourceType == kHalf ) && !supports_half)) {
|
|
if (outSourceBuffer)
|
|
free(outSourceBuffer);
|
|
return error;
|
|
}
|
|
|
|
if (!outSourceBuffer)
|
|
return 0;
|
|
|
|
// If actual source type was half, convert to float for validation.
|
|
|
|
if ( sourceType == kHalf )
|
|
validationType = kFloat;
|
|
else
|
|
validationType = sourceType;
|
|
|
|
BufferOwningPtr<char> validationSource;
|
|
|
|
if ( clFormat.image_channel_data_type == CL_UNORM_INT_101010 )
|
|
{
|
|
validationSource.reset( outSourceBuffer );
|
|
}
|
|
else
|
|
{
|
|
validationSource.reset( convert_to_expected( outSourceBuffer,
|
|
width * height * depth, sourceType, validationType, get_channel_order_channel_count(clFormat.image_channel_order) ) );
|
|
free(outSourceBuffer);
|
|
}
|
|
|
|
log_info( "- Write for %s [%4ld x %4ld x %4ld] : GL Texture : %s : %s : %s =>"
|
|
" CL Image : %s : %s \n",
|
|
GetGLTargetName(target),
|
|
width, height, depth,
|
|
GetGLFormatName( format ),
|
|
GetGLFormatName( internalFormat ),
|
|
GetGLTypeName( glType),
|
|
GetChannelOrderName( clFormat.image_channel_order ),
|
|
GetChannelTypeName( clFormat.image_channel_data_type ));
|
|
|
|
// Read the results from the GL texture.
|
|
|
|
ExplicitType readType = type;
|
|
BufferOwningPtr<char> glResults( ReadGLTexture(
|
|
target, glTexture, glBuf, width, format,
|
|
internalFormat, glType, readType, /* unused */ 1, 1 ) );
|
|
if( glResults == NULL )
|
|
return -1;
|
|
|
|
// We have to convert our input buffer to the returned type, so we can validate.
|
|
BufferOwningPtr<char> convertedGLResults;
|
|
if ( clFormat.image_channel_data_type != CL_UNORM_INT_101010 )
|
|
{
|
|
convertedGLResults.reset( convert_to_expected(
|
|
glResults, width * height * depth, readType, validationType, get_channel_order_channel_count(clFormat.image_channel_order), glType ));
|
|
}
|
|
|
|
// Validate.
|
|
|
|
int valid = 0;
|
|
if (convertedGLResults) {
|
|
if( sourceType == kFloat || sourceType == kHalf )
|
|
{
|
|
if ( clFormat.image_channel_data_type == CL_UNORM_INT_101010 )
|
|
{
|
|
valid = validate_float_results_rgb_101010( validationSource, glResults, width, height, depth, 1 );
|
|
}
|
|
else
|
|
{
|
|
valid = validate_float_results( validationSource, convertedGLResults,
|
|
width, height, depth, 1, get_channel_order_channel_count(clFormat.image_channel_order) );
|
|
}
|
|
}
|
|
else
|
|
{
|
|
valid = validate_integer_results( validationSource, convertedGLResults,
|
|
width, height, depth, 1, get_explicit_type_size( readType ) );
|
|
}
|
|
}
|
|
|
|
return valid;
|
|
}
|
|
|
|
#pragma mark -
|
|
#pragma mark Write test common entry point
|
|
|
|
// This is the main loop for all of the write tests. It iterates over the
|
|
// given formats & targets, testing a variety of sizes against each
|
|
// combination.
|
|
|
|
int test_images_write_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 err = 0;
|
|
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;
|
|
}
|
|
|
|
// Get the value of CL_DEVICE_MAX_MEM_ALLOC_SIZE
|
|
cl_ulong max_individual_allocation_size = 0;
|
|
err = clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE,
|
|
sizeof(max_individual_allocation_size),
|
|
&max_individual_allocation_size, NULL);
|
|
if (err) {
|
|
log_error("ERROR: clGetDeviceInfo failed for CL_DEVICE_MAX_MEM_ALLOC_SIZE.\n");
|
|
error++;
|
|
return error;
|
|
}
|
|
|
|
size_t total_allocation_size;
|
|
size_t fidx, tidx, sidx;
|
|
|
|
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 )
|
|
continue; // skip
|
|
}
|
|
|
|
if (formats[ fidx ].datatype == GL_UNSIGNED_INT_24_8)
|
|
{
|
|
//check if a implementation supports writing to the depth stencil formats
|
|
cl_image_format imageFormat = { CL_DEPTH_STENCIL, CL_UNORM_INT24 };
|
|
if (!is_image_format_supported(context, CL_MEM_WRITE_ONLY, (targets[tidx] == GL_TEXTURE_2D || targets[tidx] == GL_TEXTURE_RECTANGLE) ? CL_MEM_OBJECT_IMAGE2D: CL_MEM_OBJECT_IMAGE2D_ARRAY, &imageFormat))
|
|
continue;
|
|
}
|
|
|
|
if (formats[ fidx ].datatype == GL_FLOAT_32_UNSIGNED_INT_24_8_REV)
|
|
{
|
|
//check if a implementation supports writing to the depth stencil formats
|
|
cl_image_format imageFormat = { CL_DEPTH_STENCIL, CL_FLOAT};
|
|
if (!is_image_format_supported(context, CL_MEM_WRITE_ONLY, (targets[tidx] == GL_TEXTURE_2D || targets[tidx] == GL_TEXTURE_RECTANGLE) ? CL_MEM_OBJECT_IMAGE2D: CL_MEM_OBJECT_IMAGE2D_ARRAY, &imageFormat))
|
|
continue;
|
|
}
|
|
|
|
if (targets[tidx] != GL_TEXTURE_BUFFER)
|
|
log_info( "Testing image write 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 write for GL format %s : %s\n",
|
|
GetGLTargetName( targets[ tidx ] ),
|
|
GetGLFormatName( formats[ fidx ].internal ));
|
|
|
|
|
|
for (sidx = 0; sidx < nsizes; sidx++) {
|
|
|
|
// All tested formats are 4-channel formats
|
|
total_allocation_size =
|
|
sizes[sidx].width * sizes[sidx].height * sizes[sidx].depth *
|
|
4 * get_explicit_type_size( formats[ fidx ].type );
|
|
|
|
if (total_allocation_size > max_individual_allocation_size) {
|
|
log_info( "The requested allocation size (%gMB) is larger than the "
|
|
"maximum individual allocation size (%gMB)\n",
|
|
total_allocation_size/(1024.0*1024.0),
|
|
max_individual_allocation_size/(1024.0*1024.0));
|
|
log_info( "Skipping write test for %s : %s : %s : %s "
|
|
" and size (%ld, %ld, %ld)\n",
|
|
GetGLTargetName( targets[ tidx ] ),
|
|
GetGLFormatName( formats[ fidx ].internal ),
|
|
GetGLBaseFormatName( formats[ fidx ].formattype ),
|
|
GetGLTypeName( formats[ fidx ].datatype ),
|
|
sizes[sidx].width,
|
|
sizes[sidx].height,
|
|
sizes[sidx].depth);
|
|
continue;
|
|
}
|
|
#ifdef GL_VERSION_3_2
|
|
if (get_base_gl_target(targets[ tidx ]) == GL_TEXTURE_2D_MULTISAMPLE ||
|
|
get_base_gl_target(targets[ tidx ]) == GL_TEXTURE_2D_MULTISAMPLE_ARRAY)
|
|
{
|
|
bool supports_msaa;
|
|
error = supportsMsaa(context, &supports_msaa);
|
|
if( error != 0 ) return error;
|
|
if (!supports_msaa) return 0;
|
|
}
|
|
if (formats[ fidx ].formattype == GL_DEPTH_COMPONENT ||
|
|
formats[ fidx ].formattype == GL_DEPTH_STENCIL)
|
|
{
|
|
bool supports_depth;
|
|
error = supportsDepth(context, &supports_depth);
|
|
if( error != 0 ) return error;
|
|
if (!supports_depth) return 0;
|
|
}
|
|
#endif
|
|
|
|
if( test_image_format_write( context, queue,
|
|
sizes[sidx].width,
|
|
sizes[sidx].height,
|
|
sizes[sidx].depth,
|
|
targets[ tidx ],
|
|
formats[ fidx ].formattype,
|
|
formats[ fidx ].internal,
|
|
formats[ fidx ].datatype,
|
|
formats[ fidx ].type, seed ) )
|
|
{
|
|
log_error( "ERROR: Image write test failed for %s : %s : %s : %s "
|
|
" and size (%ld, %ld, %ld)\n\n",
|
|
GetGLTargetName( targets[ tidx ] ),
|
|
GetGLFormatName( formats[ fidx ].internal ),
|
|
GetGLBaseFormatName( formats[ fidx ].formattype ),
|
|
GetGLTypeName( formats[ fidx ].datatype ),
|
|
sizes[sidx].width,
|
|
sizes[sidx].height,
|
|
sizes[sidx].depth);
|
|
|
|
error++;
|
|
break; // Skip other sizes for this combination
|
|
}
|
|
}
|
|
|
|
// If we passed all sizes (check versus size loop count):
|
|
|
|
if (sidx == nsizes) {
|
|
log_info( "passed: Image write 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;
|
|
}
|