Skip to content

Commit

Permalink
Fix write-only image tests in kernel_image_methods (#330) (#1014)
Browse files Browse the repository at this point in the history 
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>
  • Loading branch information
@james-morrissey-arm @alenav01
james-morrissey-arm and alenav01 committed Nov 21, 2020
1 parent 993447a commit 5bd85b7
Show file tree
Hide file tree
Showing 6 changed files with 231 additions and 135 deletions.
61 changes: 37 additions & 24 deletions test_conformance/images/kernel_image_methods/test_1D.cpp
View file
Expand Up @@ -27,24 +27,28 @@ struct image_kernel_data
};

static const char *methodTest1DImageKernelPattern =
"typedef struct {\n"
" int width;\n"
" int channelType;\n"
" int channelOrder;\n"
" int expectedChannelType;\n"
" int expectedChannelOrder;\n"
" } image_kernel_data;\n"
"__kernel void sample_kernel( read_only image1d_t input, __global image_kernel_data *outData )\n"
"{\n"
" outData->width = get_image_width( input );\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 int test_get_1Dimage_info_single( cl_context context, cl_command_queue queue, image_descriptor *imageInfo, MTdata d )
"typedef struct {\n"
" int width;\n"
" int channelType;\n"
" int channelOrder;\n"
" int expectedChannelType;\n"
" int expectedChannelOrder;\n"
" } image_kernel_data;\n"
"__kernel void sample_kernel( %s image1d_t input, __global "
"image_kernel_data *outData )\n"
"{\n"
" outData->width = get_image_width( input );\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 int test_get_1Dimage_info_single(cl_context context,
cl_command_queue queue,
image_descriptor *imageInfo, MTdata d,
cl_mem_flags flags)
{
int error = 0;

Expand All @@ -62,7 +66,9 @@ static int test_get_1Dimage_info_single( cl_context context, cl_command_queue qu
// Construct testing source
if( gDebugTrace )
log_info( " - Creating 1D image %d ...\n", (int)imageInfo->width );
image = create_image_1d( context, (cl_mem_flags)(CL_MEM_READ_ONLY), imageInfo->format, imageInfo->width, 0, NULL, NULL, &error );

image = create_image_1d(context, flags, imageInfo->format, imageInfo->width,
0, NULL, NULL, &error);
if( image == NULL )
{
log_error( "ERROR: Unable to create 1D image of size %d (%s)", (int)imageInfo->width, IGetErrorString( error ) );
Expand All @@ -74,6 +80,8 @@ static int test_get_1Dimage_info_single( cl_context context, cl_command_queue qu

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";

if(channelTypeName && strlen(channelTypeName))
sprintf(channelTypeConstantString, "CLK_%s", &channelTypeName[3]); // replace CL_* with CLK_*
Expand All @@ -82,7 +90,7 @@ static int test_get_1Dimage_info_single( cl_context context, cl_command_queue qu
sprintf(channelOrderConstantString, "CLK_%s", &channelOrderName[3]); // replace CL_* with CLK_*

// Create a program to run against
sprintf( programSrc, methodTest1DImageKernelPattern,
sprintf(programSrc, methodTest1DImageKernelPattern, image_access_qualifier,
channelTypeConstantString, channelOrderConstantString);

//log_info("-----------------------------------\n%s\n", programSrc);
Expand Down Expand Up @@ -141,7 +149,9 @@ static int test_get_1Dimage_info_single( cl_context context, cl_command_queue qu
return error;
}

int test_get_image_info_1D( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format )
int test_get_image_info_1D(cl_device_id device, cl_context context,
cl_command_queue queue, cl_image_format *format,
cl_mem_flags flags)
{
size_t maxWidth;
cl_ulong maxAllocSize, memSize;
Expand Down Expand Up @@ -171,7 +181,8 @@ int test_get_image_info_1D( cl_device_id device, cl_context context, cl_command_
if( gDebugTrace )
log_info( " at size %d\n", (int)imageInfo.width );

int ret = test_get_1Dimage_info_single( context, queue, &imageInfo, seed );
int ret = test_get_1Dimage_info_single(context, queue, &imageInfo,
seed, flags);
if( ret )
return -1;
}
Expand All @@ -192,7 +203,8 @@ int test_get_image_info_1D( cl_device_id device, cl_context context, cl_command_
log_info( "Testing %d\n", (int)sizes[ idx ][ 0 ]);
if( gDebugTrace )
log_info( " at max size %d\n", (int)sizes[ idx ][ 0 ] );
if( test_get_1Dimage_info_single( context, queue, &imageInfo, seed ) )
if (test_get_1Dimage_info_single(context, queue, &imageInfo, seed,
flags))
return -1;
}
}
Expand Down Expand Up @@ -221,7 +233,8 @@ int test_get_image_info_1D( cl_device_id device, cl_context context, cl_command_

if( gDebugTrace )
log_info( " at size %d (row pitch %d) out of %d\n", (int)imageInfo.width, (int)imageInfo.rowPitch, (int)maxWidth );
int ret = test_get_1Dimage_info_single( context, queue, &imageInfo, seed );
int ret = test_get_1Dimage_info_single(context, queue, &imageInfo,
seed, flags);
if( ret )
return -1;
}
Expand Down
65 changes: 39 additions & 26 deletions test_conformance/images/kernel_image_methods/test_1D_array.cpp
View file
Expand Up @@ -28,26 +28,30 @@ struct image_kernel_data
};

static const char *methodTestKernelPattern =
"typedef struct {\n"
" int width;\n"
" int arraySize;\n"
" int channelType;\n"
" int channelOrder;\n"
" int expectedChannelType;\n"
" int expectedChannelOrder;\n"
" } image_kernel_data;\n"
"__kernel void sample_kernel( read_only image1d_array_t input, __global image_kernel_data *outData )\n"
"{\n"
" outData->width = get_image_width( input );\n"
" outData->arraySize = get_image_array_size( input );\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"
"}";

int test_get_1Dimage_array_info_single( cl_context context, cl_command_queue queue, image_descriptor *imageInfo, MTdata d )
"typedef struct {\n"
" int width;\n"
" int arraySize;\n"
" int channelType;\n"
" int channelOrder;\n"
" int expectedChannelType;\n"
" int expectedChannelOrder;\n"
" } image_kernel_data;\n"
"__kernel void sample_kernel( %s image1d_array_t input, __global "
"image_kernel_data *outData )\n"
"{\n"
" outData->width = get_image_width( input );\n"
" outData->arraySize = get_image_array_size( input );\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"
"}";

int test_get_1Dimage_array_info_single(cl_context context,
cl_command_queue queue,
image_descriptor *imageInfo, MTdata d,
cl_mem_flags flags)
{
int error = 0;

Expand All @@ -66,7 +70,9 @@ int test_get_1Dimage_array_info_single( cl_context context, cl_command_queue que
if( gDebugTrace )
log_info( " - Creating 1D image array %d by %d...\n", (int)imageInfo->width, (int)imageInfo->arraySize );

image = create_image_1d_array( context, (cl_mem_flags)(CL_MEM_READ_ONLY), imageInfo->format, imageInfo->width, imageInfo->arraySize, 0, 0, NULL, &error );
image = create_image_1d_array(context, flags, imageInfo->format,
imageInfo->width, imageInfo->arraySize, 0, 0,
NULL, &error);
if( image == NULL )
{
log_error( "ERROR: Unable to create 1D image array of size %d x %d (%s)", (int)imageInfo->width, (int)imageInfo->arraySize, IGetErrorString( error ) );
Expand All @@ -78,6 +84,8 @@ int test_get_1Dimage_array_info_single( cl_context context, cl_command_queue que

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";

if(channelTypeName && strlen(channelTypeName))
sprintf(channelTypeConstantString, "CLK_%s", &channelTypeName[3]); // replace CL_* with CLK_*
Expand All @@ -86,7 +94,7 @@ int test_get_1Dimage_array_info_single( cl_context context, cl_command_queue que
sprintf(channelOrderConstantString, "CLK_%s", &channelOrderName[3]); // replace CL_* with CLK_*

// Create a program to run against
sprintf( programSrc, methodTestKernelPattern,
sprintf(programSrc, methodTestKernelPattern, image_access_qualifier,
channelTypeConstantString, channelOrderConstantString);

//log_info("-----------------------------------\n%s\n", programSrc);
Expand Down Expand Up @@ -150,7 +158,9 @@ int test_get_1Dimage_array_info_single( cl_context context, cl_command_queue que
return error;
}

int test_get_image_info_1D_array( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format )
int test_get_image_info_1D_array(cl_device_id device, cl_context context,
cl_command_queue queue,
cl_image_format *format, cl_mem_flags flags)
{
size_t maxWidth, maxArraySize;
cl_ulong maxAllocSize, memSize;
Expand Down Expand Up @@ -184,7 +194,8 @@ int test_get_image_info_1D_array( cl_device_id device, cl_context context, cl_co
if( gDebugTrace )
log_info( " at size %d,%d\n", (int)imageInfo.width, (int)imageInfo.arraySize );

int ret = test_get_1Dimage_array_info_single( context, queue, &imageInfo, seed );
int ret = test_get_1Dimage_array_info_single(
context, queue, &imageInfo, seed, flags);
if( ret )
return -1;
}
Expand All @@ -208,7 +219,8 @@ int test_get_image_info_1D_array( cl_device_id device, cl_context context, cl_co
log_info( "Testing %d x %d\n", (int)sizes[ idx ][ 0 ], (int)sizes[ idx ][ 2 ]);
if( gDebugTrace )
log_info( " at max size %d,%d\n", (int)sizes[ idx ][ 0 ], (int)sizes[ idx ][ 2 ] );
if( test_get_1Dimage_array_info_single( context, queue, &imageInfo, seed ) )
if (test_get_1Dimage_array_info_single(context, queue, &imageInfo,
seed, flags))
return -1;
}
}
Expand Down Expand Up @@ -240,7 +252,8 @@ int test_get_image_info_1D_array( cl_device_id device, cl_context context, cl_co

if( gDebugTrace )
log_info( " at size %d,%d (row pitch %d) out of %d,%d\n", (int)imageInfo.width, (int)imageInfo.arraySize, (int)imageInfo.rowPitch, (int)maxWidth, (int)maxArraySize );
int ret = test_get_1Dimage_array_info_single( context, queue, &imageInfo, seed );
int ret = test_get_1Dimage_array_info_single(
context, queue, &imageInfo, seed, flags);
if( ret )
return -1;
}
Expand Down
100 changes: 60 additions & 40 deletions test_conformance/images/kernel_image_methods/test_2D.cpp
View file
Expand Up @@ -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;

Expand All @@ -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 ) );
Expand All @@ -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_*
Expand All @@ -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);
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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;
}
Expand All @@ -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;
}
}
Expand Down Expand Up @@ -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;
}
Expand Down

0 comments on commit 5bd85b7

Please sign in to comment.