ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-22 15:19:02 +00:00
Code Packages Projects Releases Wiki Activity

Fix write-only image tests in kernel_image_methods (#330) (#1014)

Browse Source 
Pass the right flags to the tests that create images to ensure valid
usage.

Fixes issue #330

Signed-off-by: James Morrissey <james.morrissey@arm.com>

Co-authored-by: Alessandro Navone <alessandro.navone@arm.com>
This commit is contained in:
james-morrissey-arm
2020-11-21 08:33:04 +00:00
committed by GitHub
parent 993447a74a
commit 5bd85b7384
6 changed files with 227 additions and 131 deletions
Download Patch File Download Diff File Expand all files Collapse all files

100
test_conformance/images/kernel_image_methods/test_2D.cpp
View File

@@ -32,38 +32,42 @@ struct image_kernel_data
};
static const char *methodTestKernelPattern =
"typedef struct {\n"
" int width;\n"
" int height;\n"
" int depth;\n"
" int widthDim;\n"
" int heightDim;\n"
" int depthDim;\n"
" int channelType;\n"
" int channelOrder;\n"
" int expectedChannelType;\n"
" int expectedChannelOrder;\n"
" } image_kernel_data;\n"
"__kernel void sample_kernel( read_only image%dd%s_t input, __global image_kernel_data *outData )\n"
"{\n"
" outData->width = get_image_width( input );\n"
" outData->height = get_image_height( input );\n"
"%s\n"
" int%d dim = get_image_dim( input );\n"
" outData->widthDim = dim.x;\n"
" outData->heightDim = dim.y;\n"
"%s\n"
" outData->channelType = get_image_channel_data_type( input );\n"
" outData->channelOrder = get_image_channel_order( input );\n"
"\n"
" outData->expectedChannelType = %s;\n"
" outData->expectedChannelOrder = %s;\n"
"}";
"typedef struct {\n"
" int width;\n"
" int height;\n"
" int depth;\n"
" int widthDim;\n"
" int heightDim;\n"
" int depthDim;\n"
" int channelType;\n"
" int channelOrder;\n"
" int expectedChannelType;\n"
" int expectedChannelOrder;\n"
" } image_kernel_data;\n"
" %s\n"
"__kernel void sample_kernel( %s image%dd%s_t input, __global "
"image_kernel_data *outData )\n"
"{\n"
" outData->width = get_image_width( input );\n"
" outData->height = get_image_height( input );\n"
"%s\n"
" int%d dim = get_image_dim( input );\n"
" outData->widthDim = dim.x;\n"
" outData->heightDim = dim.y;\n"
"%s\n"
" outData->channelType = get_image_channel_data_type( input );\n"
" outData->channelOrder = get_image_channel_order( input );\n"
"\n"
" outData->expectedChannelType = %s;\n"
" outData->expectedChannelOrder = %s;\n"
"}";
static const char *depthKernelLine = " outData->depth = get_image_depth( input );\n";
static const char *depthDimKernelLine = " outData->depthDim = dim.z;\n";
int test_get_image_info_single( cl_context context, cl_command_queue queue, image_descriptor *imageInfo, MTdata d )
int test_get_image_info_single(cl_context context, cl_command_queue queue,
image_descriptor *imageInfo, MTdata d,
cl_mem_flags flags)
{
int error = 0;
@@ -83,9 +87,13 @@ int test_get_image_info_single( cl_context context, cl_command_queue queue, imag
log_info( " - Creating image %d by %d...\n", (int)imageInfo->width, (int)imageInfo->height );
if( imageInfo->depth != 0 )
image = create_image_3d( context, (cl_mem_flags)(CL_MEM_READ_ONLY), imageInfo->format, imageInfo->width, imageInfo->height, imageInfo->depth, 0, 0, NULL, &error );
image = create_image_3d(context, flags, imageInfo->format,
imageInfo->width, imageInfo->height,
imageInfo->depth, 0, 0, NULL, &error);
else
image = create_image_2d( context, (cl_mem_flags)(CL_MEM_READ_ONLY), imageInfo->format, imageInfo->width, imageInfo->height, 0, NULL, &error );
image =
create_image_2d(context, flags, imageInfo->format, imageInfo->width,
imageInfo->height, 0, NULL, &error);
if( image == NULL )
{
log_error( "ERROR: Unable to create image of size %d x %d x %d (%s)", (int)imageInfo->width, (int)imageInfo->height, (int)imageInfo->depth, IGetErrorString( error ) );
@@ -97,6 +105,12 @@ int test_get_image_info_single( cl_context context, cl_command_queue queue, imag
const char* channelTypeName = GetChannelTypeName( imageInfo->format->image_channel_data_type );
const char* channelOrderName = GetChannelOrderName( imageInfo->format->image_channel_order );
const char *image_access_qualifier =
(flags == CL_MEM_READ_ONLY) ? "read_only" : "write_only";
const char *cl_khr_3d_image_writes_enabler = "";
if ((flags != CL_MEM_READ_ONLY) && (imageInfo->depth != 0))
cl_khr_3d_image_writes_enabler =
"#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable";
if(channelTypeName && strlen(channelTypeName))
sprintf(channelTypeConstantString, "CLK_%s", &channelTypeName[3]); // replace CL_* with CLK_*
@@ -105,12 +119,13 @@ int test_get_image_info_single( cl_context context, cl_command_queue queue, imag
sprintf(channelOrderConstantString, "CLK_%s", &channelOrderName[3]); // replace CL_* with CLK_*
// Create a program to run against
sprintf( programSrc, methodTestKernelPattern,
( imageInfo->depth != 0 ) ? 3 : 2,
(imageInfo->format->image_channel_order == CL_DEPTH) ? "_depth" : "",
( imageInfo->depth != 0 ) ? depthKernelLine : "",
( imageInfo->depth != 0 ) ? 4 : 2,
( imageInfo->depth != 0 ) ? depthDimKernelLine : "",
sprintf(programSrc, methodTestKernelPattern, cl_khr_3d_image_writes_enabler,
image_access_qualifier, (imageInfo->depth != 0) ? 3 : 2,
(imageInfo->format->image_channel_order == CL_DEPTH) ? "_depth"
: "",
(imageInfo->depth != 0) ? depthKernelLine : "",
(imageInfo->depth != 0) ? 4 : 2,
(imageInfo->depth != 0) ? depthDimKernelLine : "",
channelTypeConstantString, channelOrderConstantString);
//log_info("-----------------------------------\n%s\n", programSrc);
@@ -194,7 +209,9 @@ int test_get_image_info_single( cl_context context, cl_command_queue queue, imag
return error;
}
int test_get_image_info_2D( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format )
int test_get_image_info_2D(cl_device_id device, cl_context context,
cl_command_queue queue, cl_image_format *format,
cl_mem_flags flags)
{
size_t maxWidth, maxHeight;
cl_ulong maxAllocSize, memSize;
@@ -227,7 +244,8 @@ int test_get_image_info_2D( cl_device_id device, cl_context context, cl_command_
if( gDebugTrace )
log_info( " at size %d,%d\n", (int)imageInfo.width, (int)imageInfo.height );
int ret = test_get_image_info_single( context, queue, &imageInfo, seed );
int ret = test_get_image_info_single(context, queue, &imageInfo,
seed, flags);
if( ret )
return -1;
}
@@ -250,7 +268,8 @@ int test_get_image_info_2D( cl_device_id device, cl_context context, cl_command_
log_info( "Testing %d x %d\n", (int)sizes[ idx ][ 0 ], (int)sizes[ idx ][ 1 ]);
if( gDebugTrace )
log_info( " at max size %d,%d\n", (int)sizes[ idx ][ 0 ], (int)sizes[ idx ][ 1 ] );
if( test_get_image_info_single( context, queue, &imageInfo, seed ) )
if (test_get_image_info_single(context, queue, &imageInfo, seed,
flags))
return -1;
}
}
@@ -280,7 +299,8 @@ int test_get_image_info_2D( cl_device_id device, cl_context context, cl_command_
if( gDebugTrace )
log_info( " at size %d,%d (row pitch %d) out of %d,%d\n", (int)imageInfo.width, (int)imageInfo.height, (int)imageInfo.rowPitch, (int)maxWidth, (int)maxHeight );
int ret = test_get_image_info_single( context, queue, &imageInfo, seed );
int ret = test_get_image_info_single(context, queue, &imageInfo,
seed, flags);
if( ret )
return -1;
}