From 5bd85b7384f7a61b5a768f7d4c3dd73175412af9 Mon Sep 17 00:00:00 2001 From: james-morrissey-arm Date: Sat, 21 Nov 2020 08:33:04 +0000 Subject: [PATCH] Fix write-only image tests in kernel_image_methods (#330) (#1014) Pass the right flags to the tests that create images to ensure valid usage. Fixes issue #330 Signed-off-by: James Morrissey Co-authored-by: Alessandro Navone --- .../images/kernel_image_methods/test_1D.cpp | 59 +++++++---- .../kernel_image_methods/test_1D_array.cpp | 63 ++++++----- .../images/kernel_image_methods/test_2D.cpp | 100 +++++++++++------- .../kernel_image_methods/test_2D_array.cpp | 71 ++++++++----- .../images/kernel_image_methods/test_3D.cpp | 28 ++++- .../kernel_image_methods/test_loops.cpp | 37 +++++-- 6 files changed, 227 insertions(+), 131 deletions(-) diff --git a/test_conformance/images/kernel_image_methods/test_1D.cpp b/test_conformance/images/kernel_image_methods/test_1D.cpp index 1093eafe..1ea8eb88 100644 --- a/test_conformance/images/kernel_image_methods/test_1D.cpp +++ b/test_conformance/images/kernel_image_methods/test_1D.cpp @@ -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" -"}"; + "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 ) +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; @@ -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 ) ); @@ -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_* @@ -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); @@ -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; @@ -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; } @@ -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; } } @@ -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; } diff --git a/test_conformance/images/kernel_image_methods/test_1D_array.cpp b/test_conformance/images/kernel_image_methods/test_1D_array.cpp index 2cc39a72..18c190bb 100644 --- a/test_conformance/images/kernel_image_methods/test_1D_array.cpp +++ b/test_conformance/images/kernel_image_methods/test_1D_array.cpp @@ -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" -"}"; + "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 ) +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; @@ -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 ) ); @@ -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_* @@ -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); @@ -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; @@ -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; } @@ -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; } } @@ -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; } diff --git a/test_conformance/images/kernel_image_methods/test_2D.cpp b/test_conformance/images/kernel_image_methods/test_2D.cpp index d036c3ea..2ebc5460 100644 --- a/test_conformance/images/kernel_image_methods/test_2D.cpp +++ b/test_conformance/images/kernel_image_methods/test_2D.cpp @@ -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; } diff --git a/test_conformance/images/kernel_image_methods/test_2D_array.cpp b/test_conformance/images/kernel_image_methods/test_2D_array.cpp index 73c6db2f..98c11062 100644 --- a/test_conformance/images/kernel_image_methods/test_2D_array.cpp +++ b/test_conformance/images/kernel_image_methods/test_2D_array.cpp @@ -29,28 +29,32 @@ struct image_kernel_data }; static const char *methodTestKernelPattern = -"typedef struct {\n" -" int width;\n" -" int height;\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 %s input, __global image_kernel_data *outData )\n" -"{\n" -" outData->width = get_image_width( input );\n" -" outData->height = get_image_height( 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" -"}"; + "typedef struct {\n" + " int width;\n" + " int height;\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 %s input, __global image_kernel_data " + "*outData )\n" + "{\n" + " outData->width = get_image_width( input );\n" + " outData->height = get_image_height( 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_2Dimage_array_info_single( cl_context context, cl_command_queue queue, image_descriptor *imageInfo, MTdata d ) +int test_get_2Dimage_array_info_single(cl_context context, + cl_command_queue queue, + image_descriptor *imageInfo, MTdata d, + cl_mem_flags flags) { int error = 0; @@ -69,7 +73,9 @@ int test_get_2Dimage_array_info_single( cl_context context, cl_command_queue que if( gDebugTrace ) log_info( " - Creating 2D image array %d by %d by %d...\n", (int)imageInfo->width, (int)imageInfo->height, (int)imageInfo->arraySize ); - image = create_image_2d_array( context, (cl_mem_flags)(CL_MEM_READ_ONLY), imageInfo->format, imageInfo->width, imageInfo->height, imageInfo->arraySize, 0, 0, NULL, &error ); + image = create_image_2d_array(context, flags, imageInfo->format, + imageInfo->width, imageInfo->height, + imageInfo->arraySize, 0, 0, NULL, &error); if( image == NULL ) { log_error( "ERROR: Unable to create 2D image array of size %d x %d x %d (%s)", (int)imageInfo->width, (int)imageInfo->height, (int)imageInfo->arraySize, IGetErrorString( error ) ); @@ -81,6 +87,8 @@ int test_get_2Dimage_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_* @@ -89,8 +97,10 @@ int test_get_2Dimage_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, - (imageInfo->format->image_channel_order == CL_DEPTH) ? "image2d_array_depth_t" : "image2d_array_t" , + sprintf(programSrc, methodTestKernelPattern, image_access_qualifier, + (imageInfo->format->image_channel_order == CL_DEPTH) + ? "image2d_array_depth_t" + : "image2d_array_t", channelTypeConstantString, channelOrderConstantString); //log_info("-----------------------------------\n%s\n", programSrc); @@ -159,7 +169,9 @@ int test_get_2Dimage_array_info_single( cl_context context, cl_command_queue que return error; } -int test_get_image_info_2D_array( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format ) +int test_get_image_info_2D_array(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_image_format *format, cl_mem_flags flags) { size_t maxWidth, maxHeight, maxArraySize; cl_ulong maxAllocSize, memSize; @@ -195,7 +207,8 @@ int test_get_image_info_2D_array( cl_device_id device, cl_context context, cl_co { if( gDebugTrace ) log_info( " at size %d,%d,%d\n", (int)imageInfo.width, (int)imageInfo.height, (int)imageInfo.arraySize ); - int ret = test_get_2Dimage_array_info_single( context, queue, &imageInfo, seed ); + int ret = test_get_2Dimage_array_info_single( + context, queue, &imageInfo, seed, flags); if( ret ) return -1; } @@ -221,7 +234,8 @@ int test_get_image_info_2D_array( cl_device_id device, cl_context context, cl_co log_info( "Testing %d x %d x %d\n", (int)sizes[ idx ][ 0 ], (int)sizes[ idx ][ 1 ], (int)sizes[ idx ][ 2 ] ); if( gDebugTrace ) log_info( " at max size %d,%d,%d\n", (int)sizes[ idx ][ 0 ], (int)sizes[ idx ][ 1 ], (int)sizes[ idx ][ 2 ] ); - if( test_get_2Dimage_array_info_single( context, queue, &imageInfo, seed ) ) + if (test_get_2Dimage_array_info_single(context, queue, &imageInfo, + seed, flags)) return -1; } } @@ -257,7 +271,8 @@ int test_get_image_info_2D_array( cl_device_id device, cl_context context, cl_co if( gDebugTrace ) log_info( " at size %d,%d,%d (pitch %d,%d) out of %d,%d,%d\n", (int)imageInfo.width, (int)imageInfo.height, (int)imageInfo.arraySize, (int)imageInfo.rowPitch, (int)imageInfo.slicePitch, (int)maxWidth, (int)maxHeight, (int)maxArraySize ); - int ret = test_get_2Dimage_array_info_single( context, queue, &imageInfo, seed ); + int ret = test_get_2Dimage_array_info_single( + context, queue, &imageInfo, seed, flags); if( ret ) return -1; } diff --git a/test_conformance/images/kernel_image_methods/test_3D.cpp b/test_conformance/images/kernel_image_methods/test_3D.cpp index 3ac7a250..287005a5 100644 --- a/test_conformance/images/kernel_image_methods/test_3D.cpp +++ b/test_conformance/images/kernel_image_methods/test_3D.cpp @@ -15,9 +15,14 @@ // #include "../testBase.h" -extern int test_get_image_info_single( cl_context context, cl_command_queue queue, image_descriptor *imageInfo, MTdata d ); +extern int test_get_image_info_single(cl_context context, + cl_command_queue queue, + image_descriptor *imageInfo, MTdata d, + cl_mem_flags flags); -int test_get_image_info_3D( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format ) +int test_get_image_info_3D(cl_device_id device, cl_context context, + cl_command_queue queue, cl_image_format *format, + cl_mem_flags flags) { size_t maxWidth, maxHeight, maxDepth; cl_ulong maxAllocSize, memSize; @@ -25,6 +30,16 @@ int test_get_image_info_3D( cl_device_id device, cl_context context, cl_command_ RandomSeed seed( gRandomSeed ); size_t pixelSize; + if ((flags != CL_MEM_READ_ONLY) + && !is_extension_available(device, "cl_khr_3d_image_writes")) + { + log_info("-----------------------------------------------------\n"); + log_info("This device does not support cl_khr_3d_image_writes.\n" + "Skipping 3d image write test.\n"); + log_info("-----------------------------------------------------\n\n"); + return 0; + } + imageInfo.type = CL_MEM_OBJECT_IMAGE3D; imageInfo.format = format; pixelSize = get_pixel_size( imageInfo.format ); @@ -53,7 +68,8 @@ int test_get_image_info_3D( cl_device_id device, cl_context context, cl_command_ { if( gDebugTrace ) log_info( " at size %d,%d,%d\n", (int)imageInfo.width, (int)imageInfo.height, (int)imageInfo.depth ); - 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; } @@ -79,7 +95,8 @@ int test_get_image_info_3D( cl_device_id device, cl_context context, cl_command_ log_info( "Testing %d x %d x %d\n", (int)sizes[ idx ][ 0 ], (int)sizes[ idx ][ 1 ], (int)sizes[ idx ][ 2 ] ); if( gDebugTrace ) log_info( " at max size %d,%d,%d\n", (int)sizes[ idx ][ 0 ], (int)sizes[ idx ][ 1 ], (int)sizes[ idx ][ 2 ] ); - if( test_get_image_info_single( context, queue, &imageInfo, seed ) ) + if (test_get_image_info_single(context, queue, &imageInfo, seed, + flags)) return -1; } } @@ -115,7 +132,8 @@ int test_get_image_info_3D( cl_device_id device, cl_context context, cl_command_ if( gDebugTrace ) log_info( " at size %d,%d,%d (pitch %d,%d) out of %d,%d,%d\n", (int)imageInfo.width, (int)imageInfo.height, (int)imageInfo.depth, (int)imageInfo.rowPitch, (int)imageInfo.slicePitch, (int)maxWidth, (int)maxHeight, (int)maxDepth ); - 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; } diff --git a/test_conformance/images/kernel_image_methods/test_loops.cpp b/test_conformance/images/kernel_image_methods/test_loops.cpp index 8c0a4ead..8dfebd2f 100644 --- a/test_conformance/images/kernel_image_methods/test_loops.cpp +++ b/test_conformance/images/kernel_image_methods/test_loops.cpp @@ -18,11 +18,23 @@ extern bool gDeviceLt20; -extern int test_get_image_info_1D( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format ); -extern int test_get_image_info_2D( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format ); -extern int test_get_image_info_3D( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format ); -extern int test_get_image_info_1D_array( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format ); -extern int test_get_image_info_2D_array( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format ); +extern 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); +extern 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); +extern int test_get_image_info_3D(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_image_format *format, cl_mem_flags flags); +extern 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); +extern int test_get_image_info_2D_array(cl_device_id device, cl_context context, + cl_command_queue queue, + cl_image_format *format, + cl_mem_flags flags); int test_image_type( cl_device_id device, cl_context context, cl_command_queue queue, cl_mem_object_type imageType, cl_mem_flags flags ) { @@ -64,19 +76,24 @@ int test_image_type( cl_device_id device, cl_context context, cl_command_queue q switch (imageType) { case CL_MEM_OBJECT_IMAGE1D: - test_return = test_get_image_info_1D( device, context, queue, &formatList[ i ] ); + test_return = test_get_image_info_1D(device, context, queue, + &formatList[i], flags); break; case CL_MEM_OBJECT_IMAGE2D: - test_return = test_get_image_info_2D( device, context, queue, &formatList[ i ] ); + test_return = test_get_image_info_2D(device, context, queue, + &formatList[i], flags); break; case CL_MEM_OBJECT_IMAGE3D: - test_return = test_get_image_info_3D( device, context, queue, &formatList[ i ] ); + test_return = test_get_image_info_3D(device, context, queue, + &formatList[i], flags); break; case CL_MEM_OBJECT_IMAGE1D_ARRAY: - test_return = test_get_image_info_1D_array( device, context, queue, &formatList[ i ] ); + test_return = test_get_image_info_1D_array( + device, context, queue, &formatList[i], flags); break; case CL_MEM_OBJECT_IMAGE2D_ARRAY: - test_return = test_get_image_info_2D_array( device, context, queue, &formatList[ i ] ); + test_return = test_get_image_info_2D_array( + device, context, queue, &formatList[i], flags); break; }