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 <james.morrissey@arm.com>

Co-authored-by: Alessandro Navone <alessandro.navone@arm.com>
diff --git a/test_conformance/images/kernel_image_methods/test_1D.cpp b/test_conformance/images/kernel_image_methods/test_1D.cpp
index 1093eaf..1ea8eb8 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 @@
 };
 
 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 @@
     // 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 @@
 
     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 @@
         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 @@
     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 @@
             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 @@
             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 @@
 
             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 2cc39a7..18c190b 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 @@
 };
 
 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 @@
     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 @@
 
     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 @@
         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 @@
     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 @@
                 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 @@
             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 @@
 
             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 d036c3e..2ebc546 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 @@
 };
 
 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 @@
         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 @@
 
     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 @@
         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 @@
     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 @@
                 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 @@
             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 @@
 
             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 73c6db2..98c1106 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 @@
 };
 
 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 @@
     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 @@
 
     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 @@
         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 @@
     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 @@
                 {
                     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 @@
             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 @@
 
             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 3ac7a25..287005a 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 @@
     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 @@
                 {
                     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 @@
             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 @@
 
             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 8c0a4ea..8dfebd2 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 @@
 
         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;
         }