Enable mipmap extension pragmas (#1349)

* Enable mipmap pragmas where appopriate. * clang-format changes.
2026-03-19 06:09:01 +00:00 · 2022-05-17 16:52:40 +01:00
parent 6e6249fb48
commit d54954c7cf
10 changed files with 296 additions and 225 deletions
--- a/test_conformance/images/kernel_read_write/test_iterations.cpp
+++ b/test_conformance/images/kernel_read_write/test_iterations.cpp
@@ -39,24 +39,28 @@ static size_t reduceImageSizeRange(size_t maxDimSize) {
 }

 const char *read2DKernelSourcePattern =
-"__kernel void sample_kernel( read_only %s input,%s __global float *xOffsets, __global float *yOffsets, __global %s%s *results %s)\n"
-"{\n"
-"%s"
-"   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
-"%s"
-"%s"
-"   results[offset] = read_image%s( input, imageSampler, coords %s);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( read_only %s input,%s __global float "
+    "*xOffsets, __global float *yOffsets, __global %s%s *results %s)\n"
+    "{\n"
+    "%s"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
+    "%s"
+    "%s"
+    "   results[offset] = read_image%s( input, imageSampler, coords %s);\n"
+    "}";

 const char *read_write2DKernelSourcePattern =
-"__kernel void sample_kernel( read_write %s input,%s __global float *xOffsets, __global float *yOffsets, __global %s%s *results %s)\n"
-"{\n"
-"%s"
-"   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
-"%s"
-"%s"
-"   results[offset] = read_image%s( input, coords %s);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( read_write %s input,%s __global float "
+    "*xOffsets, __global float *yOffsets, __global %s%s *results %s)\n"
+    "{\n"
+    "%s"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
+    "%s"
+    "%s"
+    "   results[offset] = read_image%s( input, coords %s);\n"
+    "}";

 const char *intCoordKernelSource =
 "   int2 coords = (int2)( xOffsets[offset], yOffsets[offset]);\n";
@@ -1691,16 +1695,18 @@ int test_read_image_set_2D(cl_device_id device, cl_context context,
    }


-    sprintf( programSrc, KernelSourcePattern,
-            (format->image_channel_order == CL_DEPTH) ? "image2d_depth_t" : "image2d_t",
-            samplerArg, get_explicit_type_name( outputType ),
+    sprintf(programSrc, KernelSourcePattern,
+            gTestMipmaps
+                ? "#pragma OPENCL EXTENSION cl_khr_mipmap_image: enable"
+                : "",
+            (format->image_channel_order == CL_DEPTH) ? "image2d_depth_t"
+                                                      : "image2d_t",
+            samplerArg, get_explicit_type_name(outputType),
            (format->image_channel_order == CL_DEPTH) ? "" : "4",
-            gTestMipmaps?", float lod":" ",
-            samplerVar,
-            gTestMipmaps? lodOffsetSource : offsetSource,
-            floatCoords ? floatKernelSource : intCoordKernelSource,
-            readFormat,
-            gTestMipmaps?", lod":" ");
+            gTestMipmaps ? ", float lod" : " ", samplerVar,
+            gTestMipmaps ? lodOffsetSource : offsetSource,
+            floatCoords ? floatKernelSource : intCoordKernelSource, readFormat,
+            gTestMipmaps ? ", lod" : " ");

    ptr = programSrc;
    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
--- a/test_conformance/images/kernel_read_write/test_read_1D.cpp
+++ b/test_conformance/images/kernel_read_write/test_read_1D.cpp
@@ -26,24 +26,28 @@
 #endif

 const char *read1DKernelSourcePattern =
-"__kernel void sample_kernel( read_only image1d_t input,%s __global float *xOffsets, __global %s4 *results %s)\n"
-"{\n"
-"%s"
-"   int tidX = get_global_id(0);\n"
-"   int offset = tidX;\n"
-"%s"
-"   results[offset] = read_image%s( input, imageSampler, coord %s);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( read_only image1d_t input,%s __global float "
+    "*xOffsets, __global %s4 *results %s)\n"
+    "{\n"
+    "%s"
+    "   int tidX = get_global_id(0);\n"
+    "   int offset = tidX;\n"
+    "%s"
+    "   results[offset] = read_image%s( input, imageSampler, coord %s);\n"
+    "}";

 const char *read_write1DKernelSourcePattern =
-"__kernel void sample_kernel( read_write image1d_t input,%s __global float *xOffsets, __global %s4 *results %s)\n"
-"{\n"
-"%s"
-"   int tidX = get_global_id(0);\n"
-"   int offset = tidX;\n"
-"%s"
-"   results[offset] = read_image%s( input, coord %s);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( read_write image1d_t input,%s __global float "
+    "*xOffsets, __global %s4 *results %s)\n"
+    "{\n"
+    "%s"
+    "   int tidX = get_global_id(0);\n"
+    "   int offset = tidX;\n"
+    "%s"
+    "   results[offset] = read_image%s( input, coord %s);\n"
+    "}";

 const char *int1DCoordKernelSource =
 "   int coord = xOffsets[offset];\n";
@@ -1075,14 +1079,14 @@ int test_read_image_set_1D(cl_device_id device, cl_context context,
    {
        KernelSourcePattern = read1DKernelSourcePattern;
    }
-    sprintf( programSrc,
-            KernelSourcePattern,
-            samplerArg, get_explicit_type_name( outputType ),
-            gTestMipmaps ? ", float lod" : "",
-            samplerVar,
+    sprintf(programSrc, KernelSourcePattern,
+            gTestMipmaps
+                ? "#pragma OPENCL EXTENSION cl_khr_mipmap_image: enable"
+                : "",
+            samplerArg, get_explicit_type_name(outputType),
+            gTestMipmaps ? ", float lod" : "", samplerVar,
            floatCoords ? float1DKernelSource : int1DCoordKernelSource,
-            readFormat,
-            gTestMipmaps ? ", lod" : "" );
+            readFormat, gTestMipmaps ? ", lod" : "");

    ptr = programSrc;

--- a/test_conformance/images/kernel_read_write/test_read_1D_array.cpp
+++ b/test_conformance/images/kernel_read_write/test_read_1D_array.cpp
@@ -25,24 +25,28 @@
 #endif

 const char *read1DArrayKernelSourcePattern =
-"__kernel void sample_kernel( read_only image1d_array_t input,%s __global float *xOffsets, __global float *yOffsets, __global %s4 *results %s)\n"
-"{\n"
-"%s"
-"   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
-"%s"
-"%s"
-"   results[offset] = read_image%s( input, imageSampler, coords %s);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( read_only image1d_array_t input,%s __global "
+    "float *xOffsets, __global float *yOffsets, __global %s4 *results %s)\n"
+    "{\n"
+    "%s"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
+    "%s"
+    "%s"
+    "   results[offset] = read_image%s( input, imageSampler, coords %s);\n"
+    "}";

 const char *read_write1DArrayKernelSourcePattern =
-"__kernel void sample_kernel( read_write image1d_array_t input,%s __global float *xOffsets, __global float *yOffsets, __global %s4 *results %s )\n"
-"{\n"
-"%s"
-"   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
-"%s"
-"%s"
-"   results[offset] = read_image%s( input, coords %s);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( read_write image1d_array_t input,%s __global "
+    "float *xOffsets, __global float *yOffsets, __global %s4 *results %s )\n"
+    "{\n"
+    "%s"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
+    "%s"
+    "%s"
+    "   results[offset] = read_image%s( input, coords %s);\n"
+    "}";

 const char *offset1DArrayKernelSource =
 "   int offset = tidY*get_image_width(input) + tidX;\n";
@@ -1180,15 +1184,15 @@ int test_read_image_set_1D_array(cl_device_id device, cl_context context,
        KernelSourcePattern = read_write1DArrayKernelSourcePattern;
    }

-    sprintf( programSrc,
-            KernelSourcePattern,
-            samplerArg, get_explicit_type_name( outputType ),
-            gTestMipmaps ? ", float lod" : "",
-            samplerVar,
-            gTestMipmaps ? offset1DArrayLodKernelSource : offset1DArrayKernelSource,
-            floatCoords ? floatKernelSource1DArray : intCoordKernelSource1DArray,
-            readFormat,
-            gTestMipmaps ? ", lod" : "" );
+    sprintf(
+        programSrc, KernelSourcePattern,
+        gTestMipmaps ? "#pragma OPENCL EXTENSION cl_khr_mipmap_image: enable"
+                     : "",
+        samplerArg, get_explicit_type_name(outputType),
+        gTestMipmaps ? ", float lod" : "", samplerVar,
+        gTestMipmaps ? offset1DArrayLodKernelSource : offset1DArrayKernelSource,
+        floatCoords ? floatKernelSource1DArray : intCoordKernelSource1DArray,
+        readFormat, gTestMipmaps ? ", lod" : "");

    ptr = programSrc;
    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
--- a/test_conformance/images/kernel_read_write/test_read_2D_array.cpp
+++ b/test_conformance/images/kernel_read_write/test_read_2D_array.cpp
@@ -41,24 +41,32 @@ static size_t reduceImageDepth(size_t maxDepth) {
 }

 const char *read2DArrayKernelSourcePattern =
-"__kernel void sample_kernel( read_only %s input,%s __global float *xOffsets, __global float *yOffsets, __global float *zOffsets,  __global %s%s *results %s )\n"
-"{\n"
-"%s"
-"   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = get_global_id(2);\n"
-"%s"
-"%s"
-"   results[offset] = read_image%s( input, imageSampler, coords %s);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( read_only %s input,%s __global float "
+    "*xOffsets, __global float *yOffsets, __global float *zOffsets,  __global "
+    "%s%s *results %s )\n"
+    "{\n"
+    "%s"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = "
+    "get_global_id(2);\n"
+    "%s"
+    "%s"
+    "   results[offset] = read_image%s( input, imageSampler, coords %s);\n"
+    "}";

 const char *read_write2DArrayKernelSourcePattern =
-"__kernel void sample_kernel( read_write %s input,%s __global float *xOffsets, __global float *yOffsets, __global float *zOffsets,  __global %s%s *results %s)\n"
-"{\n"
-"%s"
-"   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = get_global_id(2);\n"
-"%s"
-"%s"
-"   results[offset] = read_image%s( input, coords %s);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( read_write %s input,%s __global float "
+    "*xOffsets, __global float *yOffsets, __global float *zOffsets,  __global "
+    "%s%s *results %s)\n"
+    "{\n"
+    "%s"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = "
+    "get_global_id(2);\n"
+    "%s"
+    "%s"
+    "   results[offset] = read_image%s( input, coords %s);\n"
+    "}";

 const char* offset2DarraySource ="   int offset = tidZ*get_image_width(input)*get_image_height(input) + tidY*get_image_width(input) + tidX;\n";
 const char* offset2DarraySourceLod =
@@ -1412,17 +1420,16 @@ int test_read_image_set_2D_array(cl_device_id device, cl_context context,
    }

    // Construct the source
-    sprintf( programSrc,
-            KernelSourcePattern,
-            imageType,
-            samplerArg, get_explicit_type_name( outputType ),
-            imageElement,
-            gTestMipmaps ? ", float lod" : " ",
-            samplerVar,
+    sprintf(programSrc, KernelSourcePattern,
+            gTestMipmaps
+                ? "#pragma OPENCL EXTENSION cl_khr_mipmap_image: enable"
+                : "",
+            imageType, samplerArg, get_explicit_type_name(outputType),
+            imageElement, gTestMipmaps ? ", float lod" : " ", samplerVar,
            gTestMipmaps ? offset2DarraySourceLod : offset2DarraySource,
-            floatCoords ? float2DArrayUnnormalizedCoordKernelSource : int2DArrayCoordKernelSource,
-            readFormat,
-            gTestMipmaps ? ", lod" : " " );
+            floatCoords ? float2DArrayUnnormalizedCoordKernelSource
+                        : int2DArrayCoordKernelSource,
+            readFormat, gTestMipmaps ? ", lod" : " ");

    ptr = programSrc;
    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
--- a/test_conformance/images/kernel_read_write/test_read_3D.cpp
+++ b/test_conformance/images/kernel_read_write/test_read_3D.cpp
@@ -36,24 +36,32 @@ static size_t reduceImageDepth(size_t maxDimSize, RandomSeed& seed) {


 const char *read3DKernelSourcePattern =
-"__kernel void sample_kernel( read_only image3d_t input,%s __global float *xOffsets, __global float *yOffsets, __global float *zOffsets,  __global %s4 *results %s)\n"
-"{\n"
-"%s"
-"   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = get_global_id(2);\n"
-"%s"
-"%s"
-"   results[offset] = read_image%s( input, imageSampler, coords %s);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( read_only image3d_t input,%s __global float "
+    "*xOffsets, __global float *yOffsets, __global float *zOffsets,  __global "
+    "%s4 *results %s)\n"
+    "{\n"
+    "%s"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = "
+    "get_global_id(2);\n"
+    "%s"
+    "%s"
+    "   results[offset] = read_image%s( input, imageSampler, coords %s);\n"
+    "}";

 const char *read_write3DKernelSourcePattern =
-"__kernel void sample_kernel( read_write image3d_t input,%s __global float *xOffsets, __global float *yOffsets, __global float *zOffsets,  __global %s4 *results %s)\n"
-"{\n"
-"%s"
-"   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = get_global_id(2);\n"
-"%s"
-"%s"
-"   results[offset] = read_image%s( input, coords %s);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( read_write image3d_t input,%s __global float "
+    "*xOffsets, __global float *yOffsets, __global float *zOffsets,  __global "
+    "%s4 *results %s)\n"
+    "{\n"
+    "%s"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = "
+    "get_global_id(2);\n"
+    "%s"
+    "%s"
+    "   results[offset] = read_image%s( input, coords %s);\n"
+    "}";

 const char *offset3DKernelSource =
 "   int offset = tidZ*get_image_width(input)*get_image_height(input) + tidY*get_image_width(input) + tidX;\n";
@@ -137,15 +145,16 @@ int test_read_image_set_3D(cl_device_id device, cl_context context,
        KernelSourcePattern = read_write3DKernelSourcePattern;
    }

-    sprintf( programSrc,
-            KernelSourcePattern,
-            samplerArg, get_explicit_type_name( outputType ),
-            gTestMipmaps? ", float lod": " ",
-            samplerVar,
-            gTestMipmaps? offset3DLodKernelSource: offset3DKernelSource,
-            floatCoords ? float3DUnnormalizedCoordKernelSource : int3DCoordKernelSource,
-            readFormat,
-            gTestMipmaps? ",lod":" ");
+    sprintf(programSrc, KernelSourcePattern,
+            gTestMipmaps
+                ? "#pragma OPENCL EXTENSION cl_khr_mipmap_image: enable"
+                : "",
+            samplerArg, get_explicit_type_name(outputType),
+            gTestMipmaps ? ", float lod" : " ", samplerVar,
+            gTestMipmaps ? offset3DLodKernelSource : offset3DKernelSource,
+            floatCoords ? float3DUnnormalizedCoordKernelSource
+                        : int3DCoordKernelSource,
+            readFormat, gTestMipmaps ? ",lod" : " ");

    ptr = programSrc;
    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
--- a/test_conformance/images/kernel_read_write/test_write_1D.cpp
+++ b/test_conformance/images/kernel_read_write/test_write_1D.cpp
@@ -27,20 +27,24 @@ extern bool validate_float_write_results( float *expected, float *actual, image_
 extern bool validate_half_write_results( cl_half *expected, cl_half *actual, image_descriptor* imageInfo );

 const char *readwrite1DKernelSourcePattern =
-"__kernel void sample_kernel( __global %s4 *input, read_write image1d_t output %s)\n"
-"{\n"
-"   int tidX = get_global_id(0);\n"
-"   int offset = tidX;\n"
-"   write_image%s( output, tidX %s, input[ offset ]);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( __global %s4 *input, read_write image1d_t "
+    "output %s)\n"
+    "{\n"
+    "   int tidX = get_global_id(0);\n"
+    "   int offset = tidX;\n"
+    "   write_image%s( output, tidX %s, input[ offset ]);\n"
+    "}";

 const char *write1DKernelSourcePattern =
-"__kernel void sample_kernel( __global %s4 *input, write_only image1d_t output %s)\n"
-"{\n"
-"   int tidX = get_global_id(0);\n"
-"   int offset = tidX;\n"
-"   write_image%s( output, tidX %s, input[ offset ]);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( __global %s4 *input, write_only image1d_t "
+    "output %s)\n"
+    "{\n"
+    "   int tidX = get_global_id(0);\n"
+    "   int offset = tidX;\n"
+    "   write_image%s( output, tidX %s, input[ offset ]);\n"
+    "}";

 int test_write_image_1D( cl_device_id device, cl_context context, cl_command_queue queue, cl_kernel kernel,
                     image_descriptor *imageInfo, ExplicitType inputType, MTdata d )
@@ -614,12 +618,14 @@ int test_write_image_1D_set(cl_device_id device, cl_context context,
        KernelSourcePattern = readwrite1DKernelSourcePattern;
    }

-    sprintf( programSrc,
-             KernelSourcePattern,
-             get_explicit_type_name( inputType ),
-             gTestMipmaps ? ", int lod" : "",
-             readFormat,
-             gTestMipmaps ? ", lod" :"" );
+    sprintf(
+        programSrc, KernelSourcePattern,
+        gTestMipmaps
+            ? "#pragma OPENCL EXTENSION cl_khr_mipmap_image: enable\n#pragma "
+              "OPENCL EXTENSION cl_khr_mipmap_image_writes: enable"
+            : "",
+        get_explicit_type_name(inputType), gTestMipmaps ? ", int lod" : "",
+        readFormat, gTestMipmaps ? ", lod" : "");

    ptr = programSrc;
    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
--- a/test_conformance/images/kernel_read_write/test_write_1D_array.cpp
+++ b/test_conformance/images/kernel_read_write/test_write_1D_array.cpp
@@ -27,20 +27,24 @@ extern bool validate_float_write_results( float *expected, float *actual, image_
 extern bool validate_half_write_results( cl_half *expected, cl_half *actual, image_descriptor *imageInfo );

 const char *readwrite1DArrayKernelSourcePattern =
-"__kernel void sample_kernel( __global %s4 *input, read_write image1d_array_t output %s)\n"
-"{\n"
-"   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
-"%s"
-"   write_image%s( output, (int2)( tidX, tidY )%s, input[ offset ]);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( __global %s4 *input, read_write "
+    "image1d_array_t output %s)\n"
+    "{\n"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
+    "%s"
+    "   write_image%s( output, (int2)( tidX, tidY )%s, input[ offset ]);\n"
+    "}";

 const char *write1DArrayKernelSourcePattern =
-"__kernel void sample_kernel( __global %s4 *input, write_only image1d_array_t output %s)\n"
-"{\n"
-"   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
-"%s"
-"   write_image%s( output, (int2)( tidX, tidY ) %s, input[ offset ]);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( __global %s4 *input, write_only "
+    "image1d_array_t output %s)\n"
+    "{\n"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
+    "%s"
+    "   write_image%s( output, (int2)( tidX, tidY ) %s, input[ offset ]);\n"
+    "}";

 const char *offset1DArraySource =
 "   int offset = tidY*get_image_width(output) + tidX;\n";
@@ -637,13 +641,15 @@ int test_write_image_1D_array_set(cl_device_id device, cl_context context,
    }
    // Construct the source
    // Construct the source
-    sprintf( programSrc,
-             KernelSourcePattern,
-             get_explicit_type_name( inputType ),
-             gTestMipmaps ? ", int lod" : "",
-             gTestMipmaps ? offset1DArrayLodSource : offset1DArraySource,
-             readFormat,
-             gTestMipmaps ? ", lod" :"" );
+    sprintf(
+        programSrc, KernelSourcePattern,
+        gTestMipmaps
+            ? "#pragma OPENCL EXTENSION cl_khr_mipmap_image: enable\n#pragma "
+              "OPENCL EXTENSION cl_khr_mipmap_image_writes: enable"
+            : "",
+        get_explicit_type_name(inputType), gTestMipmaps ? ", int lod" : "",
+        gTestMipmaps ? offset1DArrayLodSource : offset1DArraySource, readFormat,
+        gTestMipmaps ? ", lod" : "");

    ptr = programSrc;
    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
--- a/test_conformance/images/kernel_read_write/test_write_2D_array.cpp
+++ b/test_conformance/images/kernel_read_write/test_write_2D_array.cpp
@@ -49,20 +49,28 @@ static size_t reduceImageDepth(size_t maxDepth) {
 }

 const char *write2DArrayKernelSourcePattern =
-"__kernel void sample_kernel( __global %s%s *input, write_only %s output %s)\n"
-"{\n"
-"   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = get_global_id(2);\n"
-"%s"
-"   write_image%s( output, (int4)( tidX, tidY, tidZ, 0 ) %s, input[ offset ]);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( __global %s%s *input, write_only %s output "
+    "%s)\n"
+    "{\n"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = "
+    "get_global_id(2);\n"
+    "%s"
+    "   write_image%s( output, (int4)( tidX, tidY, tidZ, 0 ) %s, input[ offset "
+    "]);\n"
+    "}";

 const char *readwrite2DArrayKernelSourcePattern =
-"__kernel void sample_kernel( __global %s%s *input, read_write %s output %s)\n"
-"{\n"
-"   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = get_global_id(2);\n"
-"%s"
-"   write_image%s( output, (int4)( tidX, tidY, tidZ, 0 ) %s, input[ offset ] );\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( __global %s%s *input, read_write %s output "
+    "%s)\n"
+    "{\n"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = "
+    "get_global_id(2);\n"
+    "%s"
+    "   write_image%s( output, (int4)( tidX, tidY, tidZ, 0 ) %s, input[ offset "
+    "] );\n"
+    "}";

 const char *offset2DArrayKernelSource =
 "   int offset = tidZ*get_image_width(output)*get_image_height(output) + tidY*get_image_width(output) + tidX;\n";
@@ -671,15 +679,19 @@ int test_write_image_2D_array_set(cl_device_id device, cl_context context,
    }
    // Construct the source
    // Construct the source
-    sprintf( programSrc,
-             KernelSourcePattern,
-             get_explicit_type_name( inputType ),
-             (format->image_channel_order == CL_DEPTH) ? "" : "4",
-             (format->image_channel_order == CL_DEPTH) ? "image2d_array_depth_t" : "image2d_array_t",
-             gTestMipmaps ? " , int lod" : "",
-             gTestMipmaps ? offset2DArrayLodKernelSource : offset2DArrayKernelSource,
-             readFormat,
-             gTestMipmaps ? ", lod" : "" );
+    sprintf(
+        programSrc, KernelSourcePattern,
+        gTestMipmaps
+            ? "#pragma OPENCL EXTENSION cl_khr_mipmap_image: enable\n#pragma "
+              "OPENCL EXTENSION cl_khr_mipmap_image_writes: enable"
+            : "",
+        get_explicit_type_name(inputType),
+        (format->image_channel_order == CL_DEPTH) ? "" : "4",
+        (format->image_channel_order == CL_DEPTH) ? "image2d_array_depth_t"
+                                                  : "image2d_array_t",
+        gTestMipmaps ? " , int lod" : "",
+        gTestMipmaps ? offset2DArrayLodKernelSource : offset2DArrayKernelSource,
+        readFormat, gTestMipmaps ? ", lod" : "");

    ptr = programSrc;
    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
--- a/test_conformance/images/kernel_read_write/test_write_3D.cpp
+++ b/test_conformance/images/kernel_read_write/test_write_3D.cpp
@@ -46,22 +46,30 @@ static size_t reduceImageDepth(size_t maxDimSize, MTdata& seed) {


 const char *write3DKernelSourcePattern =
-"%s"
-"__kernel void sample_kernel( __global %s4 *input, write_only image3d_t output %s )\n"
-"{\n"
-"   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = get_global_id(2);\n"
-"%s"
-"   write_image%s( output, (int4)( tidX, tidY, tidZ, 0 ) %s, input[ offset ]);\n"
-"}";
+    "%s"
+    "%s\n"
+    "__kernel void sample_kernel( __global %s4 *input, write_only image3d_t "
+    "output %s )\n"
+    "{\n"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = "
+    "get_global_id(2);\n"
+    "%s"
+    "   write_image%s( output, (int4)( tidX, tidY, tidZ, 0 ) %s, input[ offset "
+    "]);\n"
+    "}";

 const char *readwrite3DKernelSourcePattern =
-"%s"
-"__kernel void sample_kernel( __global %s4 *input, read_write image3d_t output %s )\n"
-"{\n"
-"   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = get_global_id(2);\n"
-"%s"
-"   write_image%s( output, (int4)( tidX, tidY, tidZ, 0 ) %s, input[ offset ]);\n"
-"}";
+    "%s"
+    "%s\n"
+    "__kernel void sample_kernel( __global %s4 *input, read_write image3d_t "
+    "output %s )\n"
+    "{\n"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1), tidZ = "
+    "get_global_id(2);\n"
+    "%s"
+    "   write_image%s( output, (int4)( tidX, tidY, tidZ, 0 ) %s, input[ offset "
+    "]);\n"
+    "}";

 const char *khr3DWritesPragma =
 "#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable\n";
@@ -678,14 +686,15 @@ int test_write_image_3D_set(cl_device_id device, cl_context context,
    }

    // Construct the source
-    sprintf( programSrc,
-             KernelSourcePattern,
-             gTestMipmaps ? "" : khr3DWritesPragma,
-             get_explicit_type_name( inputType ),
-             gTestMipmaps ? ", int lod" : "",
-             gTestMipmaps ? offset3DLodSource : offset3DSource,
-             readFormat,
-             gTestMipmaps ? ", lod" : "" );
+    sprintf(
+        programSrc, KernelSourcePattern, khr3DWritesPragma,
+        gTestMipmaps
+            ? "#pragma OPENCL EXTENSION cl_khr_mipmap_image: enable\n#pragma "
+              "OPENCL EXTENSION cl_khr_mipmap_image_writes: enable"
+            : "",
+        get_explicit_type_name(inputType), gTestMipmaps ? ", int lod" : "",
+        gTestMipmaps ? offset3DLodSource : offset3DSource, readFormat,
+        gTestMipmaps ? ", lod" : "");

    ptr = programSrc;
    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
--- a/test_conformance/images/kernel_read_write/test_write_image.cpp
+++ b/test_conformance/images/kernel_read_write/test_write_image.cpp
@@ -47,20 +47,24 @@ extern bool validate_float_write_results( float *expected, float *actual, image_
 extern bool validate_half_write_results( cl_half *expected, cl_half *actual, image_descriptor *imageInfo );

 const char *writeKernelSourcePattern =
-"__kernel void sample_kernel( __global %s%s *input, write_only %s output %s)\n"
-"{\n"
-"   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
-"%s"
-"   write_image%s( output, (int2)( tidX, tidY ) %s, input[ offset ]);\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( __global %s%s *input, write_only %s output "
+    "%s)\n"
+    "{\n"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
+    "%s"
+    "   write_image%s( output, (int2)( tidX, tidY ) %s, input[ offset ]);\n"
+    "}";

 const char *read_writeKernelSourcePattern =
-"__kernel void sample_kernel( __global %s%s *input, read_write %s output %s)\n"
-"{\n"
-"   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
-"%s"
-"   write_image%s( output, (int2)( tidX, tidY )%s, input[ offset ] );\n"
-"}";
+    "%s\n"
+    "__kernel void sample_kernel( __global %s%s *input, read_write %s output "
+    "%s)\n"
+    "{\n"
+    "   int tidX = get_global_id(0), tidY = get_global_id(1);\n"
+    "%s"
+    "   write_image%s( output, (int2)( tidX, tidY )%s, input[ offset ] );\n"
+    "}";

 const char *offset2DKernelSource =
 "   int offset = tidY*get_image_width(output) + tidX;\n";
@@ -728,15 +732,19 @@ int test_write_image_set(cl_device_id device, cl_context context,
    }

    // Construct the source
-    sprintf( programSrc,
-             KernelSourcePattern,
-             get_explicit_type_name( inputType ),
-             (format->image_channel_order == CL_DEPTH) ? "" : "4",
-             (format->image_channel_order == CL_DEPTH) ? "image2d_depth_t" : "image2d_t",
-             gTestMipmaps ? ", int lod" : "",
-             gTestMipmaps ? offset2DLodKernelSource : offset2DKernelSource,
-             readFormat,
-             gTestMipmaps ? ", lod" : "" );
+    sprintf(
+        programSrc, KernelSourcePattern,
+        gTestMipmaps
+            ? "#pragma OPENCL EXTENSION cl_khr_mipmap_image: enable\n#pragma "
+              "OPENCL EXTENSION cl_khr_mipmap_image_writes: enable"
+            : "",
+        get_explicit_type_name(inputType),
+        (format->image_channel_order == CL_DEPTH) ? "" : "4",
+        (format->image_channel_order == CL_DEPTH) ? "image2d_depth_t"
+                                                  : "image2d_t",
+        gTestMipmaps ? ", int lod" : "",
+        gTestMipmaps ? offset2DLodKernelSource : offset2DKernelSource,
+        readFormat, gTestMipmaps ? ", lod" : "");

    ptr = programSrc;
    error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,