From d54954c7cfd4311d12d076b205ee632b0d6cc151 Mon Sep 17 00:00:00 2001 From: Jeremy Kemp Date: Tue, 17 May 2022 16:52:40 +0100 Subject: [PATCH] Enable mipmap extension pragmas (#1349) * Enable mipmap pragmas where appopriate. * clang-format changes. --- .../kernel_read_write/test_iterations.cpp | 56 ++++++++++-------- .../images/kernel_read_write/test_read_1D.cpp | 50 ++++++++-------- .../kernel_read_write/test_read_1D_array.cpp | 54 +++++++++-------- .../kernel_read_write/test_read_2D_array.cpp | 59 +++++++++++-------- .../images/kernel_read_write/test_read_3D.cpp | 59 +++++++++++-------- .../kernel_read_write/test_write_1D.cpp | 42 +++++++------ .../kernel_read_write/test_write_1D_array.cpp | 44 ++++++++------ .../kernel_read_write/test_write_2D_array.cpp | 54 ++++++++++------- .../kernel_read_write/test_write_3D.cpp | 53 ++++++++++------- .../kernel_read_write/test_write_image.cpp | 50 +++++++++------- 10 files changed, 296 insertions(+), 225 deletions(-) diff --git a/test_conformance/images/kernel_read_write/test_iterations.cpp b/test_conformance/images/kernel_read_write/test_iterations.cpp index 2f5c75a7..05aed02c 100644 --- 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, diff --git a/test_conformance/images/kernel_read_write/test_read_1D.cpp b/test_conformance/images/kernel_read_write/test_read_1D.cpp index e9306fc4..2a722088 100644 --- 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; diff --git a/test_conformance/images/kernel_read_write/test_read_1D_array.cpp b/test_conformance/images/kernel_read_write/test_read_1D_array.cpp index 2f4e4d3b..a8009420 100644 --- 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, diff --git a/test_conformance/images/kernel_read_write/test_read_2D_array.cpp b/test_conformance/images/kernel_read_write/test_read_2D_array.cpp index d71bfec4..533a0fe8 100644 --- 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, diff --git a/test_conformance/images/kernel_read_write/test_read_3D.cpp b/test_conformance/images/kernel_read_write/test_read_3D.cpp index 860114fb..cec77bf0 100644 --- 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, diff --git a/test_conformance/images/kernel_read_write/test_write_1D.cpp b/test_conformance/images/kernel_read_write/test_write_1D.cpp index 1556a76a..5f726796 100644 --- 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, diff --git a/test_conformance/images/kernel_read_write/test_write_1D_array.cpp b/test_conformance/images/kernel_read_write/test_write_1D_array.cpp index e9aa8d2a..f9024405 100644 --- 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, diff --git a/test_conformance/images/kernel_read_write/test_write_2D_array.cpp b/test_conformance/images/kernel_read_write/test_write_2D_array.cpp index 5bca7124..c1c56994 100644 --- 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, diff --git a/test_conformance/images/kernel_read_write/test_write_3D.cpp b/test_conformance/images/kernel_read_write/test_write_3D.cpp index d9a69627..9da93695 100644 --- 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, diff --git a/test_conformance/images/kernel_read_write/test_write_image.cpp b/test_conformance/images/kernel_read_write/test_write_image.cpp index 9cc9698c..29626971 100644 --- 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,