From 1c04050b5b018f2502e9723f3c24a2811e4dd0f4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?K=C3=A9vin=20Petit?= Date: Wed, 11 Dec 2019 10:36:45 +0000 Subject: [PATCH] Use old-style sampler creation in basic suite and remove duplicate compatibility tests (#510) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit As agreed in the WG calls, the query tests are sufficient to cover both old-style and new-style sampler and command queue creation entrypoints. Use old-style entrypoints in all tests that don't require 2.x features to get compatibility with 1.x implementations. This makes it possible to remove duplicate compatibility tests. Contributes to #494. Signed-off-by: Kévin Petit --- test_conformance/basic/test_image_multipass.c | 18 +- test_conformance/basic/test_image_param.c | 9 +- test_conformance/basic/test_image_r8.c | 9 +- test_conformance/basic/test_imagedim.c | 24 +- test_conformance/basic/test_imagenpot.c | 9 +- .../basic/test_multireadimagemultifmt.c | 9 +- .../basic/test_multireadimageonefmt.c | 9 +- test_conformance/basic/test_readimage3d.c | 9 +- .../basic/test_readimage3d_fp32.c | 9 +- .../basic/test_readimage3d_int16.c | 9 +- test_conformance/basic/test_readimage_fp32.c | 9 +- test_conformance/basic/test_readimage_int16.c | 9 +- .../test_conformance/basic/CMakeLists.txt | 9 +- .../test_conformance/basic/main.c | 15 - .../basic/test_image_multipass.c | 643 ------------------ .../test_conformance/basic/test_image_param.c | 251 ------- .../test_conformance/basic/test_image_r8.c | 176 ----- .../test_conformance/basic/test_imagedim.c | 514 -------------- .../test_conformance/basic/test_imagenpot.c | 220 ------ .../basic/test_multireadimagemultifmt.c | 230 ------- .../basic/test_multireadimageonefmt.c | 198 ------ .../test_conformance/basic/test_readimage3d.c | 230 ------- .../basic/test_readimage3d_fp32.c | 147 ---- .../basic/test_readimage3d_int16.c | 146 ---- .../basic/test_readimage_fp32.c | 167 ----- .../basic/test_readimage_int16.c | 166 ----- 26 files changed, 32 insertions(+), 3212 deletions(-) delete mode 100644 test_conformance/compatibility/test_conformance/basic/test_image_multipass.c delete mode 100644 test_conformance/compatibility/test_conformance/basic/test_image_param.c delete mode 100644 test_conformance/compatibility/test_conformance/basic/test_image_r8.c delete mode 100644 test_conformance/compatibility/test_conformance/basic/test_imagedim.c delete mode 100644 test_conformance/compatibility/test_conformance/basic/test_imagenpot.c delete mode 100644 test_conformance/compatibility/test_conformance/basic/test_multireadimagemultifmt.c delete mode 100644 test_conformance/compatibility/test_conformance/basic/test_multireadimageonefmt.c delete mode 100644 test_conformance/compatibility/test_conformance/basic/test_readimage3d.c delete mode 100644 test_conformance/compatibility/test_conformance/basic/test_readimage3d_fp32.c delete mode 100644 test_conformance/compatibility/test_conformance/basic/test_readimage3d_int16.c delete mode 100644 test_conformance/compatibility/test_conformance/basic/test_readimage_fp32.c delete mode 100644 test_conformance/compatibility/test_conformance/basic/test_readimage_int16.c diff --git a/test_conformance/basic/test_image_multipass.c b/test_conformance/basic/test_image_multipass.c index 5f75c936..cd91a136 100644 --- a/test_conformance/basic/test_image_multipass.c +++ b/test_conformance/basic/test_image_multipass.c @@ -286,13 +286,8 @@ test_image_multipass_integer_coord(cl_device_id device, cl_context context, cl_c clReleaseProgram(program[1]); } - cl_sampler_properties properties[] = { - CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, - CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, - CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, - 0 }; - cl_sampler sampler = clCreateSamplerWithProperties(context, properties, &err); - test_error(err, "clCreateSamplerWithProperties failed"); + cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); { size_t threads[3] = {0, 0, 0}; @@ -532,13 +527,8 @@ test_image_multipass_float_coord(cl_device_id device, cl_context context, cl_com clReleaseProgram(program[1]); } - cl_sampler_properties properties[] = { - CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, - CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, - CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, - 0 }; - cl_sampler sampler = clCreateSamplerWithProperties(context, properties, &err); - test_error(err, "clCreateSamplerWithProperties failed"); + cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); { size_t threads[3] = {0, 0, 0}; diff --git a/test_conformance/basic/test_image_param.c b/test_conformance/basic/test_image_param.c index 5d3656fc..c8fc1b4a 100644 --- a/test_conformance/basic/test_image_param.c +++ b/test_conformance/basic/test_image_param.c @@ -239,13 +239,8 @@ int test_image_param(cl_device_id device, cl_context context, cl_command_queue q test_error( error, "Unable to create testing kernel" ); // Also create a sampler to use for all the runs - cl_sampler_properties properties[] = { - CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, - CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, - CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, - 0 }; - clSamplerWrapper sampler = clCreateSamplerWithProperties(context, properties, &error); - test_error(error, "clCreateSamplerWithProperties failed"); + clSamplerWrapper sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &error ); + test_error( error, "clCreateSampler failed" ); // Set up the arguments for each and queue for( i = 0, idx = 0; i < numSizes; i++ ) diff --git a/test_conformance/basic/test_image_r8.c b/test_conformance/basic/test_image_r8.c index bf8596ca..7805c1b5 100644 --- a/test_conformance/basic/test_image_r8.c +++ b/test_conformance/basic/test_image_r8.c @@ -129,13 +129,8 @@ test_image_r8(cl_device_id device, cl_context context, cl_command_queue queue, i return -1; } - cl_sampler_properties properties[] = { - CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, - CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, - CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, - 0 }; - cl_sampler sampler = clCreateSamplerWithProperties(context, properties, &err); - test_error(err, "clCreateSamplerWithProperties failed"); + cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); diff --git a/test_conformance/basic/test_imagedim.c b/test_conformance/basic/test_imagedim.c index 8d14a7f6..6d8cdb36 100644 --- a/test_conformance/basic/test_imagedim.c +++ b/test_conformance/basic/test_imagedim.c @@ -117,13 +117,8 @@ test_imagedim_pow2(cl_device_id device, cl_context context, cl_command_queue que max_mem_size = (cl_ulong)SIZE_MAX; } - cl_sampler_properties properties[] = { - CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, - CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, - CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, - 0 }; - cl_sampler sampler = clCreateSamplerWithProperties(context, properties, &err); - test_error(err, "clCreateSamplerWithProperties failed"); + cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); max_img_width = (int)max_image2d_width; max_img_height = (int)max_image2d_height; @@ -154,9 +149,9 @@ test_imagedim_pow2(cl_device_id device, cl_context context, cl_command_queue que log_info("Adjusted maximum image size to test is %d x %d, which is a max mem size of %gMB.\n", max_img_width, max_img_height, (max_img_width*max_img_height*4)/(1024.0*1024.0)); - d = init_genrand( gRandomSeed ); - input_ptr = generate_8888_image(max_img_width, max_img_height, d); - output_ptr = (unsigned char*)malloc(sizeof(unsigned char) * 4 * max_img_width * max_img_height); + d = init_genrand( gRandomSeed ); + input_ptr = generate_8888_image(max_img_width, max_img_height, d); + output_ptr = (unsigned char*)malloc(sizeof(unsigned char) * 4 * max_img_width * max_img_height); // test power of 2 width, height starting at 1 to 4K for (i=1,i2=0; i<=max_img_height; i<<=1,i2++) @@ -328,13 +323,8 @@ test_imagedim_non_pow2(cl_device_id device, cl_context context, cl_command_queue log_info("Device reported max image sizes of %lu x %lu, and max mem size of %gMB.\n", max_image2d_width, max_image2d_height, max_mem_size/(1024.0*1024.0)); - cl_sampler_properties properties[] = { - CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, - CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, - CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, - 0 }; - cl_sampler sampler = clCreateSamplerWithProperties(context, properties, &err); - test_error(err, "clCreateSamplerWithProperties failed"); + cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); max_img_width = (int)max_image2d_width; max_img_height = (int)max_image2d_height; diff --git a/test_conformance/basic/test_imagenpot.c b/test_conformance/basic/test_imagenpot.c index 72a97675..4713c30e 100644 --- a/test_conformance/basic/test_imagenpot.c +++ b/test_conformance/basic/test_imagenpot.c @@ -150,13 +150,8 @@ test_imagenpot(cl_device_id device_id, cl_context context, cl_command_queue queu return -1; } - cl_sampler_properties properties[] = { - CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, - CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, - CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, - 0 }; - cl_sampler sampler = clCreateSamplerWithProperties(context, properties, &err); - test_error(err, "clCreateSamplerWithProperties failed"); + cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); diff --git a/test_conformance/basic/test_multireadimagemultifmt.c b/test_conformance/basic/test_multireadimagemultifmt.c index fd15d284..5c93d2fc 100644 --- a/test_conformance/basic/test_multireadimagemultifmt.c +++ b/test_conformance/basic/test_multireadimagemultifmt.c @@ -181,13 +181,8 @@ test_mri_multiple(cl_device_id device, cl_context context, cl_command_queue queu if (err) return -1; - cl_sampler_properties properties[] = { - CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, - CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, - CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, - 0 }; - cl_sampler sampler = clCreateSamplerWithProperties(context, properties, &err); - test_error(err, "clCreateSamplerWithProperties failed"); + cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); for (i=0; i<4; i++) err |= clSetKernelArg(kernel, i,sizeof streams[i], &streams[i]); diff --git a/test_conformance/basic/test_multireadimageonefmt.c b/test_conformance/basic/test_multireadimageonefmt.c index f39d8912..b37c8414 100644 --- a/test_conformance/basic/test_multireadimageonefmt.c +++ b/test_conformance/basic/test_multireadimageonefmt.c @@ -147,13 +147,8 @@ int test_mri_one(cl_device_id device, cl_context context, cl_command_queue queue if (err) return -1; - cl_sampler_properties properties[] = { - CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, - CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, - CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, - 0 }; - cl_sampler sampler = clCreateSamplerWithProperties(context, properties, &err); - test_error(err, "clCreateSamplerWithProperties failed"); + cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); err = clSetKernelArg(kernel, 0, sizeof i, &i); err |= clSetKernelArg(kernel, 1, sizeof err, &err); diff --git a/test_conformance/basic/test_readimage3d.c b/test_conformance/basic/test_readimage3d.c index f962a352..95a5b5a1 100644 --- a/test_conformance/basic/test_readimage3d.c +++ b/test_conformance/basic/test_readimage3d.c @@ -167,13 +167,8 @@ int test_readimage3d(cl_device_id device, cl_context context, cl_command_queue q if (err) return -1; - cl_sampler_properties properties[] = { - CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, - CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, - CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, - 0 }; - cl_sampler sampler = clCreateSamplerWithProperties(context, properties, &err); - test_error(err, "clCreateSamplerWithProperties failed"); + cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); err = clSetKernelArg(kernel[0], 0, sizeof streams[0], &streams[0]); err |= clSetKernelArg(kernel[0], 1, sizeof streams[2], &streams[2]); diff --git a/test_conformance/basic/test_readimage3d_fp32.c b/test_conformance/basic/test_readimage3d_fp32.c index fb0a786e..2658d365 100644 --- a/test_conformance/basic/test_readimage3d_fp32.c +++ b/test_conformance/basic/test_readimage3d_fp32.c @@ -114,13 +114,8 @@ int test_readimage3d_fp32(cl_device_id device, cl_context context, cl_command_qu if (err) return -1; - cl_sampler_properties properties[] = { - CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, - CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, - CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, - 0 }; - cl_sampler sampler = clCreateSamplerWithProperties(context, properties, &err); - test_error(err, "clCreateSamplerWithProperties failed"); + cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); diff --git a/test_conformance/basic/test_readimage3d_int16.c b/test_conformance/basic/test_readimage3d_int16.c index 6b0f7713..690d72a4 100644 --- a/test_conformance/basic/test_readimage3d_int16.c +++ b/test_conformance/basic/test_readimage3d_int16.c @@ -113,13 +113,8 @@ int test_readimage3d_int16(cl_device_id device, cl_context context, cl_command_q if (err) return -1; - cl_sampler_properties properties[] = { - CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, - CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, - CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, - 0 }; - cl_sampler sampler = clCreateSamplerWithProperties(context, properties, &err); - test_error(err, "clCreateSamplerWithProperties failed"); + cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); diff --git a/test_conformance/basic/test_readimage_fp32.c b/test_conformance/basic/test_readimage_fp32.c index 11e8369f..aa9a82f1 100644 --- a/test_conformance/basic/test_readimage_fp32.c +++ b/test_conformance/basic/test_readimage_fp32.c @@ -123,13 +123,8 @@ int test_readimage_fp32(cl_device_id device, cl_context context, cl_command_queu if (err) return -1; - cl_sampler_properties properties[] = { - CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, - CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, - CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, - 0 }; - cl_sampler sampler = clCreateSamplerWithProperties(context, properties, &err); - test_error(err, "clCreateSamplerWithProperties failed"); + cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); diff --git a/test_conformance/basic/test_readimage_int16.c b/test_conformance/basic/test_readimage_int16.c index b0e216ca..fe2912f5 100644 --- a/test_conformance/basic/test_readimage_int16.c +++ b/test_conformance/basic/test_readimage_int16.c @@ -122,13 +122,8 @@ int test_readimage_int16(cl_device_id device, cl_context context, cl_command_que if (err) return -1; - cl_sampler_properties properties[] = { - CL_SAMPLER_NORMALIZED_COORDS, CL_FALSE, - CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE, - CL_SAMPLER_FILTER_MODE, CL_FILTER_NEAREST, - 0 }; - cl_sampler sampler = clCreateSamplerWithProperties(context, properties, &err); - test_error(err, "clCreateSamplerWithProperties failed"); + cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); + test_error(err, "clCreateSampler failed"); err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); diff --git a/test_conformance/compatibility/test_conformance/basic/CMakeLists.txt b/test_conformance/compatibility/test_conformance/basic/CMakeLists.txt index 143ba737..37d48b4a 100644 --- a/test_conformance/compatibility/test_conformance/basic/CMakeLists.txt +++ b/test_conformance/compatibility/test_conformance/basic/CMakeLists.txt @@ -2,15 +2,8 @@ set(MODULE_NAME COMPATIBILITY_BASIC) set(${MODULE_NAME}_SOURCES main.c - test_readimage.c test_readimage_int16.c test_readimage_fp32.c - test_readimage3d.c test_readimage3d_int16.c test_readimage3d_fp32.c + test_readimage.c test_writeimage.c - test_multireadimageonefmt.c test_multireadimagemultifmt.c - test_imagedim.c - test_image_multipass.c - test_image_param.c - test_imagenpot.c - test_image_r8.c test_imagearraycopy3d.c test_async_copy.cpp test_sizeof.c diff --git a/test_conformance/compatibility/test_conformance/basic/main.c b/test_conformance/compatibility/test_conformance/basic/main.c index 3678992e..385ab65f 100644 --- a/test_conformance/compatibility/test_conformance/basic/main.c +++ b/test_conformance/compatibility/test_conformance/basic/main.c @@ -33,25 +33,10 @@ bool gTestRounding = false; test_definition test_list[] = { ADD_TEST( sizeof ), ADD_TEST( readimage ), - ADD_TEST( readimage_int16 ), - ADD_TEST( readimage_fp32 ), ADD_TEST( writeimage ), - ADD_TEST( mri_one ), - ADD_TEST( mri_multiple ), - ADD_TEST( image_r8 ), - ADD_TEST( readimage3d ), - ADD_TEST( readimage3d_int16 ), - ADD_TEST( readimage3d_fp32 ), ADD_TEST( bufferreadwriterect ), ADD_TEST( imagearraycopy3d ), - ADD_TEST( imagenpot ), - - ADD_TEST( imagedim_pow2 ), - ADD_TEST( imagedim_non_pow2 ), - ADD_TEST( image_param ), - ADD_TEST( image_multipass_integer_coord ), - ADD_TEST( image_multipass_float_coord ), ADD_TEST( async_copy_global_to_local ), ADD_TEST( async_copy_local_to_global ), diff --git a/test_conformance/compatibility/test_conformance/basic/test_image_multipass.c b/test_conformance/compatibility/test_conformance/basic/test_image_multipass.c deleted file mode 100644 index fbc2e506..00000000 --- a/test_conformance/compatibility/test_conformance/basic/test_image_multipass.c +++ /dev/null @@ -1,643 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "harness/compat.h" - -#include -#include -#include -#include - - -#include "procs.h" - -static const char *image_to_image_kernel_integer_coord_code = -"\n" -"__kernel void image_to_image_copy(read_only image2d_t srcimg, write_only image2d_t dstimg, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" float4 color;\n" -"\n" -" color = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y));\n" -" write_imagef(dstimg, (int2)(tid_x, tid_y), color);\n" -"\n" -"}\n"; - -static const char *image_to_image_kernel_float_coord_code = -"\n" -"__kernel void image_to_image_copy(read_only image2d_t srcimg, write_only image2d_t dstimg, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" float4 color;\n" -"\n" -" color = read_imagef(srcimg, sampler, (float2)((float)tid_x, (float)tid_y));\n" -" write_imagef(dstimg, (int2)(tid_x, tid_y), color);\n" -"\n" -"}\n"; - - -static const char *image_sum_kernel_integer_coord_code = -"\n" -"__kernel void image_sum(read_only image2d_t srcimg0, read_only image2d_t srcimg1, write_only image2d_t dstimg, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" float4 color0;\n" -" float4 color1;\n" -"\n" -" color0 = read_imagef(srcimg0, sampler, (int2)(tid_x, tid_y));\n" -" color1 = read_imagef(srcimg1, sampler, (int2)(tid_x, tid_y));\n" -" write_imagef(dstimg, (int2)(tid_x, tid_y), color0 + color1);\n" -"\n" -"}\n"; - - -static const char *image_sum_kernel_float_coord_code = -"\n" -"__kernel void image_sum(read_only image2d_t srcimg0, read_only image2d_t srcimg1, write_only image2d_t dstimg, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" float4 color0;\n" -" float4 color1;\n" -"\n" -" color0 = read_imagef(srcimg0, sampler, (float2)((float)tid_x, (float)tid_y));\n" -" color1 = read_imagef(srcimg1, sampler, (float2)((float)tid_x, (float)tid_y));\n" -" write_imagef(dstimg,(int2)(tid_x, tid_y), color0 + color1);\n" -"\n" -"}\n"; - - -static unsigned char * -generate_initial_byte_image(int w, int h, int num_elements, unsigned char value) -{ - unsigned char *ptr = (unsigned char*)malloc(w * h * num_elements); - int i; - - for (i = 0; i < w*h*num_elements; i++) - ptr[i] = value; - - return ptr; -} - -static unsigned char * -generate_expected_byte_image(unsigned char **input_data, int num_inputs, int w, int h, int num_elements) -{ - unsigned char *ptr = (unsigned char*)malloc(w * h * num_elements); - int i; - - for (i = 0; i < w*h*num_elements; i++) - { - int j; - ptr[i] = 0; - for (j = 0; j < num_inputs; j++) - { - unsigned char *input = *(input_data + j); - ptr[i] += input[i]; - } - } - - return ptr; -} - - -static unsigned char * -generate_byte_image(int w, int h, int num_elements, MTdata d) -{ - unsigned char *ptr = (unsigned char*)malloc(w * h * num_elements); - int i; - - for (i = 0; i < w*h*num_elements; i++) - ptr[i] = (unsigned char)genrand_int32(d) & 31; - - return ptr; -} - -static int -verify_byte_image(unsigned char *image, unsigned char *outptr, int w, int h, int num_elements) -{ - int i; - - for (i = 0; i < w*h*num_elements; i++) - { - if (outptr[i] != image[i]) - { - return -1; - } - } - return 0; -} - -int -test_image_multipass_integer_coord(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) -{ - int img_width = 512; - int img_height = 512; - cl_image_format img_format; - - int num_input_streams = 8; - cl_mem *input_streams; - cl_mem accum_streams[2]; - unsigned char *expected_output; - unsigned char *output_ptr; - cl_kernel kernel[2]; - int err; - - PASSIVE_REQUIRE_IMAGE_SUPPORT( device ) - - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_UNORM_INT8; - - expected_output = (unsigned char*)malloc(sizeof(unsigned char) * 4 * img_width * img_height); - output_ptr = (unsigned char*)malloc(sizeof(unsigned char) * 4 * img_width * img_height); - - // Create the accum images with initial data. - { - unsigned char *initial_data; - cl_mem_flags flags; - - initial_data = generate_initial_byte_image(img_width, img_height, 4, 0xF0); - flags = (cl_mem_flags)(CL_MEM_READ_WRITE); - - accum_streams[0] = create_image_2d(context, flags, &img_format, img_width, img_height, 0, NULL, NULL); - if (!accum_streams[0]) - { - log_error("create_image_2d failed\n"); - free(expected_output); - free(output_ptr); - return -1; - } - - size_t origin[3] = {0, 0, 0}, region[3] = {img_width, img_height, 1}; - err = clEnqueueWriteImage(queue, accum_streams[0], CL_TRUE, - origin, region, 0, 0, - initial_data, 0, NULL, NULL); - if (err) - { - log_error("clWriteImage failed: %d\n", err); - free(expected_output); - free(output_ptr); - return -1; - } - - accum_streams[1] = create_image_2d(context, flags, &img_format, img_width, img_height, 0, NULL, NULL); - if (!accum_streams[1]) - { - log_error("create_image_2d failed\n"); - free(expected_output); - free(output_ptr); - return -1; - } - err = clEnqueueWriteImage(queue, accum_streams[1], CL_TRUE, - origin, region, 0, 0, - initial_data, 0, NULL, NULL); - if (err) - { - log_error("clWriteImage failed: %d\n", err); - free(expected_output); - free(output_ptr); - return -1; - } - - free(initial_data); - } - - // Set up the input data. - { - cl_mem_flags flags; - unsigned char **input_data = (unsigned char **)malloc(sizeof(unsigned char*) * num_input_streams); - MTdata d; - - input_streams = (cl_mem*)malloc(sizeof(cl_mem) * num_input_streams); - flags = (cl_mem_flags)(CL_MEM_READ_WRITE); - - int i; - d = init_genrand( gRandomSeed ); - for ( i = 0; i < num_input_streams; i++) - { - input_data[i] = generate_byte_image(img_width, img_height, 4, d); - input_streams[i] = create_image_2d(context, flags, &img_format, img_width, img_height, 0, NULL, NULL); - if (!input_streams[i]) - { - log_error("create_image_2d failed\n"); - free_mtdata(d); - free(expected_output); - free(output_ptr); - return -1; - } - - size_t origin[3] = {0, 0, 0}, region[3] = {img_width, img_height, 1}; - err = clEnqueueWriteImage(queue, input_streams[i], CL_TRUE, - origin, region, 0, 0, - input_data[i], 0, NULL, NULL); - if (err) - { - log_error("clWriteImage failed: %d\n", err); - free_mtdata(d); - free(expected_output); - free(output_ptr); - free(input_streams); - return -1; - } - - - } - free_mtdata(d); d = NULL; - expected_output = generate_expected_byte_image(input_data, num_input_streams, img_width, img_height, 4); - for ( i = 0; i < num_input_streams; i++) - { - free(input_data[i]); - } - free( input_data ); - } - - // Set up the kernels. - { - cl_program program[4]; - - err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &image_to_image_kernel_integer_coord_code, "image_to_image_copy"); - if (err) - { - log_error("Failed to create kernel 0: %d\n", err); - return -1; - } - err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &image_sum_kernel_integer_coord_code, "image_sum"); - if (err) - { - log_error("Failed to create kernel 1: %d\n", err); - return -1; - } - clReleaseProgram(program[0]); - clReleaseProgram(program[1]); - } - - cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); - test_error(err, "clCreateSampler failed"); - - { - size_t threads[3] = {0, 0, 0}; - threads[0] = (size_t)img_width; - threads[1] = (size_t)img_height; - int i; - - { - cl_mem accum_input; - cl_mem accum_output; - - err = clSetKernelArg(kernel[0], 0, sizeof input_streams[0], &input_streams[0]); - err |= clSetKernelArg(kernel[0], 1, sizeof accum_streams[0], &accum_streams[0]); - err |= clSetKernelArg(kernel[0], 2, sizeof sampler, &sampler); - if (err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed\n"); - return -1; - } - err = clEnqueueNDRangeKernel( queue, kernel[0], 2, NULL, threads, NULL, 0, NULL, NULL ); - if (err != CL_SUCCESS) - { - log_error("clEnqueueNDRangeKernel failed\n"); - return -1; - } - - for (i = 1; i < num_input_streams; i++) - { - accum_input = accum_streams[(i-1)%2]; - accum_output = accum_streams[i%2]; - - err = clSetKernelArg(kernel[1], 0, sizeof accum_input, &accum_input); - err |= clSetKernelArg(kernel[1], 1, sizeof input_streams[i], &input_streams[i]); - err |= clSetKernelArg(kernel[1], 2, sizeof accum_output, &accum_output); - err |= clSetKernelArg(kernel[1], 3, sizeof sampler, &sampler); - - if (err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed\n"); - return -1; - } - err = clEnqueueNDRangeKernel( queue, kernel[1], 2, NULL, threads, NULL, 0, NULL, NULL ); - if (err != CL_SUCCESS) - { - log_error("clEnqueueNDRangeKernel failed\n"); - return -1; - } - } - - // Copy the last accum into the other one. - accum_input = accum_streams[(i-1)%2]; - accum_output = accum_streams[i%2]; - err = clSetKernelArg(kernel[0], 0, sizeof accum_input, &accum_input); - err |= clSetKernelArg(kernel[0], 1, sizeof accum_output, &accum_output); - if (err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed\n"); - return -1; - } - err = clEnqueueNDRangeKernel( queue, kernel[0], 2, NULL, threads, NULL, 0, NULL, NULL ); - if (err != CL_SUCCESS) - { - log_error("clEnqueueNDRangeKernel failed\n"); - return -1; - } - - size_t origin[3] = {0, 0, 0}, region[3] = {img_width, img_height, 1}; - err = clEnqueueReadImage(queue, accum_output, CL_TRUE, - origin, region, 0, 0, - (void *)output_ptr, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clReadImage failed\n"); - return -1; - } - err = verify_byte_image(expected_output, output_ptr, img_width, img_height, 4); - if (err) - { - log_error("IMAGE_MULTIPASS test failed.\n"); - } - else - { - log_info("IMAGE_MULTIPASS test passed\n"); - } - } - - clReleaseSampler(sampler); - } - - - // cleanup - clReleaseMemObject(accum_streams[0]); - clReleaseMemObject(accum_streams[1]); - { - int i; - for (i = 0; i < num_input_streams; i++) - { - clReleaseMemObject(input_streams[i]); - } - } - free(input_streams); - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - free(expected_output); - free(output_ptr); - - return err; -} - -int -test_image_multipass_float_coord(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) -{ - int img_width = 512; - int img_height = 512; - cl_image_format img_format; - - int num_input_streams = 8; - cl_mem *input_streams; - cl_mem accum_streams[2]; - unsigned char *expected_output; - unsigned char *output_ptr; - cl_kernel kernel[2]; - int err; - - PASSIVE_REQUIRE_IMAGE_SUPPORT( device ) - - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_UNORM_INT8; - - output_ptr = (unsigned char*)malloc(sizeof(unsigned char) * 4 * img_width * img_height); - - // Create the accum images with initial data. - { - unsigned char *initial_data; - cl_mem_flags flags; - - initial_data = generate_initial_byte_image(img_width, img_height, 4, 0xF0); - flags = (cl_mem_flags)(CL_MEM_READ_WRITE); - - accum_streams[0] = create_image_2d(context, flags, &img_format, img_width, img_height, 0, NULL, NULL); - if (!accum_streams[0]) - { - log_error("create_image_2d failed\n"); - return -1; - } - - size_t origin[3] = {0, 0, 0}, region[3] = {img_width, img_height, 1}; - err = clEnqueueWriteImage(queue, accum_streams[0], CL_TRUE, - origin, region, 0, 0, - initial_data, 0, NULL, NULL); - if (err) - { - log_error("clWriteImage failed: %d\n", err); - return -1; - } - - accum_streams[1] = create_image_2d(context, flags, &img_format, img_width, img_height, 0, NULL, NULL); - if (!accum_streams[1]) - { - log_error("create_image_2d failed\n"); - return -1; - } - err = clEnqueueWriteImage(queue, accum_streams[1], CL_TRUE, - origin, region, 0, 0, - initial_data, 0, NULL, NULL); - if (err) - { - log_error("clWriteImage failed: %d\n", err); - return -1; - } - - free(initial_data); - } - - // Set up the input data. - { - cl_mem_flags flags; - unsigned char **input_data = (unsigned char **)malloc(sizeof(unsigned char*) * num_input_streams); - MTdata d; - - input_streams = (cl_mem*)malloc(sizeof(cl_mem) * num_input_streams); - flags = (cl_mem_flags)(CL_MEM_READ_WRITE); - - int i; - d = init_genrand( gRandomSeed ); - for ( i = 0; i < num_input_streams; i++) - { - input_data[i] = generate_byte_image(img_width, img_height, 4, d); - input_streams[i] = create_image_2d(context, flags, &img_format, img_width, img_height, 0, NULL, NULL); - if (!input_streams[i]) - { - log_error("create_image_2d failed\n"); - free(input_data); - free(input_streams); - return -1; - } - - size_t origin[3] = {0, 0, 0}, region[3] = {img_width, img_height, 1}; - err = clEnqueueWriteImage(queue, input_streams[i], CL_TRUE, - origin, region, 0, 0, - input_data[i], 0, NULL, NULL); - if (err) - { - log_error("clWriteImage failed: %d\n", err); - free(input_data); - free(input_streams); - return -1; - } - } - free_mtdata(d); d = NULL; - expected_output = generate_expected_byte_image(input_data, num_input_streams, img_width, img_height, 4); - for ( i = 0; i < num_input_streams; i++) - { - free(input_data[i]); - } - free(input_data); - } - - // Set up the kernels. - { - cl_program program[2]; - - err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &image_to_image_kernel_float_coord_code, "image_to_image_copy"); - if (err) - { - log_error("Failed to create kernel 2: %d\n", err); - return -1; - } - err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &image_sum_kernel_float_coord_code, "image_sum"); - if (err) - { - log_error("Failed to create kernel 3: %d\n", err); - return -1; - } - - clReleaseProgram(program[0]); - clReleaseProgram(program[1]); - } - - cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); - test_error(err, "clCreateSampler failed"); - - { - size_t threads[3] = {0, 0, 0}; - threads[0] = (size_t)img_width; - threads[1] = (size_t)img_height; - int i; - - { - cl_mem accum_input; - cl_mem accum_output; - - err = clSetKernelArg(kernel[0], 0, sizeof input_streams[0], &input_streams[0]); - err |= clSetKernelArg(kernel[0], 1, sizeof accum_streams[0], &accum_streams[0]); - err |= clSetKernelArg(kernel[0], 2, sizeof sampler, &sampler); - if (err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed\n"); - return -1; - } - err = clEnqueueNDRangeKernel( queue, kernel[0], 2, NULL, threads, NULL, 0, NULL, NULL ); - if (err != CL_SUCCESS) - { - log_error("clEnqueueNDRangeKernel failed\n"); - return -1; - } - - for (i = 1; i < num_input_streams; i++) - { - accum_input = accum_streams[(i-1)%2]; - accum_output = accum_streams[i%2]; - - err = clSetKernelArg(kernel[1], 0, sizeof accum_input, &accum_input); - err |= clSetKernelArg(kernel[1], 1, sizeof input_streams[i], &input_streams[i]); - err |= clSetKernelArg(kernel[1], 2, sizeof accum_output, &accum_output); - err |= clSetKernelArg(kernel[1], 3, sizeof sampler, &sampler); - - if (err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed\n"); - return -1; - } - err = clEnqueueNDRangeKernel( queue, kernel[1], 2, NULL, threads, NULL, 0, NULL, NULL ); - if (err != CL_SUCCESS) - { - log_error("clEnqueueNDRangeKernel failed\n"); - return -1; - } - } - - // Copy the last accum into the other one. - accum_input = accum_streams[(i-1)%2]; - accum_output = accum_streams[i%2]; - err = clSetKernelArg(kernel[0], 0, sizeof accum_input, &accum_input); - err |= clSetKernelArg(kernel[0], 1, sizeof accum_output, &accum_output); - if (err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed\n"); - return -1; - } - err = clEnqueueNDRangeKernel( queue, kernel[0], 2, NULL, threads, NULL, 0, NULL, NULL ); - if (err != CL_SUCCESS) - { - log_error("clEnqueueNDRangeKernel failed\n"); - return -1; - } - - size_t origin[3] = {0, 0, 0}, region[3] = {img_width, img_height, 1}; - err = clEnqueueReadImage(queue, accum_output, CL_TRUE, - origin, region, 0, 0, - (void *)output_ptr, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clReadImage failed\n"); - return -1; - } - err = verify_byte_image(expected_output, output_ptr, img_width, img_height, 4); - if (err) - { - log_error("IMAGE_MULTIPASS test failed.\n"); - } - else - { - log_info("IMAGE_MULTIPASS test passed\n"); - } - } - - } - - - // cleanup - clReleaseSampler(sampler); - clReleaseMemObject(accum_streams[0]); - clReleaseMemObject(accum_streams[1]); - { - int i; - for (i = 0; i < num_input_streams; i++) - { - clReleaseMemObject(input_streams[i]); - } - } - clReleaseKernel(kernel[0]); - clReleaseKernel(kernel[1]); - free(expected_output); - free(output_ptr); - free(input_streams); - - return err; -} - - - - - diff --git a/test_conformance/compatibility/test_conformance/basic/test_image_param.c b/test_conformance/compatibility/test_conformance/basic/test_image_param.c deleted file mode 100644 index 229975ed..00000000 --- a/test_conformance/compatibility/test_conformance/basic/test_image_param.c +++ /dev/null @@ -1,251 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "harness/compat.h" - -#include -#include -#include -#include - - -#include "procs.h" -#include "harness/typeWrappers.h" -#include "harness/imageHelpers.h" -#include "harness/conversions.h" - - -static const char *param_kernel[] = { -"__kernel void test_fn(read_only image2d_t srcimg, sampler_t sampler, __global float4 *results )\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" results[ tid_y * get_image_width( srcimg ) + tid_x ] = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y));\n" -"\n" -"}\n" }; - -int validate_results( size_t width, size_t height, cl_image_format &format, char *inputData, cl_float *actualResults ) -{ - for( size_t i = 0; i < width * height; i++ ) - { - cl_float expected[ 4 ], tolerance; - - switch( format.image_channel_data_type ) - { - case CL_UNORM_INT8: - { - cl_uchar *p = (cl_uchar *)inputData; - expected[ 0 ] = p[ 0 ] / 255.f; - expected[ 1 ] = p[ 1 ] / 255.f; - expected[ 2 ] = p[ 2 ] / 255.f; - expected[ 3 ] = p[ 3 ] / 255.f; - tolerance = 1.f / 255.f; - break; - } - case CL_SNORM_INT8: - { - cl_char *p = (cl_char *)inputData; - expected[ 0 ] = fmaxf( p[ 0 ] / 127.f, -1.f ); - expected[ 1 ] = fmaxf( p[ 1 ] / 127.f, -1.f ); - expected[ 2 ] = fmaxf( p[ 2 ] / 127.f, -1.f ); - expected[ 3 ] = fmaxf( p[ 3 ] / 127.f, -1.f ); - tolerance = 1.f / 127.f; - break; - } - case CL_UNSIGNED_INT8: - { - cl_uchar *p = (cl_uchar *)inputData; - expected[ 0 ] = p[ 0 ]; - expected[ 1 ] = p[ 1 ]; - expected[ 2 ] = p[ 2 ]; - expected[ 3 ] = p[ 3 ]; - tolerance = 1.f / 127.f; - break; - } - case CL_SIGNED_INT8: - { - cl_short *p = (cl_short *)inputData; - expected[ 0 ] = p[ 0 ]; - expected[ 1 ] = p[ 1 ]; - expected[ 2 ] = p[ 2 ]; - expected[ 3 ] = p[ 3 ]; - tolerance = 1.f / 127.f; - break; - } - case CL_UNORM_INT16: - { - cl_ushort *p = (cl_ushort *)inputData; - expected[ 0 ] = p[ 0 ] / 65535.f; - expected[ 1 ] = p[ 1 ] / 65535.f; - expected[ 2 ] = p[ 2 ] / 65535.f; - expected[ 3 ] = p[ 3 ] / 65535.f; - tolerance = 1.f / 65535.f; - break; - } - case CL_UNSIGNED_INT32: - { - cl_uint *p = (cl_uint *)inputData; - expected[ 0 ] = p[ 0 ]; - expected[ 1 ] = p[ 1 ]; - expected[ 2 ] = p[ 2 ]; - expected[ 3 ] = p[ 3 ]; - tolerance = 0.0001f; - break; - } - case CL_FLOAT: - { - cl_float *p = (cl_float *)inputData; - expected[ 0 ] = p[ 0 ]; - expected[ 1 ] = p[ 1 ]; - expected[ 2 ] = p[ 2 ]; - expected[ 3 ] = p[ 3 ]; - tolerance = 0.0001f; - break; - } - default: - // Should never get here - break; - } - - if( format.image_channel_order == CL_BGRA ) - { - cl_float tmp = expected[ 0 ]; - expected[ 0 ] = expected[ 2 ]; - expected[ 2 ] = tmp; - } - - // Within an error tolerance, make sure the results match - cl_float error1 = fabsf( expected[ 0 ] - actualResults[ 0 ] ); - cl_float error2 = fabsf( expected[ 1 ] - actualResults[ 1 ] ); - cl_float error3 = fabsf( expected[ 2 ] - actualResults[ 2 ] ); - cl_float error4 = fabsf( expected[ 3 ] - actualResults[ 3 ] ); - - if( error1 > tolerance || error2 > tolerance || error3 > tolerance || error4 > tolerance ) - { - log_error( "ERROR: Sample %d did not validate against expected results for %d x %d %s:%s image\n", (int)i, (int)width, (int)height, - GetChannelOrderName( format.image_channel_order ), GetChannelTypeName( format.image_channel_data_type ) ); - log_error( " Expected: %f %f %f %f\n", (float)expected[ 0 ], (float)expected[ 1 ], (float)expected[ 2 ], (float)expected[ 3 ] ); - log_error( " Actual: %f %f %f %f\n", (float)actualResults[ 0 ], (float)actualResults[ 1 ], (float)actualResults[ 2 ], (float)actualResults[ 3 ] ); - - // Check real quick a special case error here - cl_float error1 = fabsf( expected[ 3 ] - actualResults[ 0 ] ); - cl_float error2 = fabsf( expected[ 2 ] - actualResults[ 1 ] ); - cl_float error3 = fabsf( expected[ 1 ] - actualResults[ 2 ] ); - cl_float error4 = fabsf( expected[ 0 ] - actualResults[ 3 ] ); - if( error1 <= tolerance && error2 <= tolerance && error3 <= tolerance && error4 <= tolerance ) - { - log_error( "\t(Kernel did not respect change in channel order)\n" ); - } - return -1; - } - - // Increment and go - actualResults += 4; - inputData += get_format_type_size( &format ) * 4; - } - - return 0; -} - -int test_image_param(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) -{ - size_t sizes[] = { 64, 100, 128, 250, 512 }; - cl_image_format formats[] = { { CL_RGBA, CL_UNORM_INT8 }, { CL_RGBA, CL_UNORM_INT16 }, { CL_RGBA, CL_FLOAT }, { CL_BGRA, CL_UNORM_INT8 } }; - ExplicitType types[] = { kUChar, kUShort, kFloat, kUChar }; - int error; - size_t i, j, idx; - size_t threads[ 2 ]; - MTdata d; - - const size_t numSizes = sizeof( sizes ) / sizeof( sizes[ 0 ] ); - const size_t numFormats = sizeof( formats ) / sizeof( formats[ 0 ] ); - const size_t numAttempts = numSizes * numFormats; - - - clProgramWrapper program; - clKernelWrapper kernel; - clMemWrapper streams[ numAttempts ][ 2 ]; - BufferOwningPtr inputs[ numAttempts ]; - - PASSIVE_REQUIRE_IMAGE_SUPPORT( device ) - - d = init_genrand( gRandomSeed ); - for( i = 0, idx = 0; i < numSizes; i++ ) - { - for( j = 0; j < numFormats; j++, idx++ ) - { - // For each attempt, we create a pair: an input image, whose parameters keep changing, and an output buffer - // that we can read values from. The output buffer will remain consistent to ensure that any changes we - // witness are due to the image changes - inputs[ idx ].reset(create_random_data( types[ j ], d, sizes[ i ] * sizes[ i ] * 4 )); - - streams[ idx ][ 0 ] = create_image_2d( context, CL_MEM_COPY_HOST_PTR, &formats[ j ], sizes[ i ], sizes[ i ], 0, inputs[ idx ], &error ); - { - char err_str[256]; - sprintf(err_str, "Unable to create input image for format %s order %s" , - GetChannelOrderName( formats[j].image_channel_order ), - GetChannelTypeName( formats[j].image_channel_data_type )); - test_error( error, err_str); - } - - streams[ idx ][ 1 ] = clCreateBuffer( context, CL_MEM_READ_WRITE, sizes[ i ] * sizes[ i ] * 4 * sizeof( cl_float ), NULL, &error ); - test_error( error, "Unable to create output buffer" ); - } - } - free_mtdata(d); d = NULL; - - // Create a single kernel to use for all the tests - error = create_single_kernel_helper( context, &program, &kernel, 1, param_kernel, "test_fn" ); - test_error( error, "Unable to create testing kernel" ); - - // Also create a sampler to use for all the runs - clSamplerWrapper sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &error ); - test_error( error, "clCreateSampler failed" ); - - // Set up the arguments for each and queue - for( i = 0, idx = 0; i < numSizes; i++ ) - { - for( j = 0; j < numFormats; j++, idx++ ) - { - error = clSetKernelArg( kernel, 0, sizeof( streams[ idx ][ 0 ] ), &streams[ idx ][ 0 ] ); - error |= clSetKernelArg( kernel, 1, sizeof( sampler ), &sampler ); - error |= clSetKernelArg( kernel, 2, sizeof( streams[ idx ][ 1 ] ), &streams[ idx ][ 1 ]); - test_error( error, "Unable to set kernel arguments" ); - - threads[ 0 ] = threads[ 1 ] = (size_t)sizes[ i ]; - - error = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, NULL, 0, NULL, NULL ); - test_error( error, "clEnqueueNDRangeKernel failed" ); - } - } - - // Now go through each combo and validate the results - for( i = 0, idx = 0; i < numSizes; i++ ) - { - for( j = 0; j < numFormats; j++, idx++ ) - { - BufferOwningPtr output(malloc(sizeof(cl_float) * sizes[ i ] * sizes[ i ] * 4 )); - - error = clEnqueueReadBuffer( queue, streams[ idx ][ 1 ], CL_TRUE, 0, sizes[ i ] * sizes[ i ] * 4 * sizeof( cl_float ), output, 0, NULL, NULL ); - test_error( error, "Unable to read results" ); - - error = validate_results( sizes[ i ], sizes[ i ], formats[ j ], inputs[ idx ], output ); - if( error ) - return -1; - } - } - - return 0; -} diff --git a/test_conformance/compatibility/test_conformance/basic/test_image_r8.c b/test_conformance/compatibility/test_conformance/basic/test_image_r8.c deleted file mode 100644 index 37f2bdba..00000000 --- a/test_conformance/compatibility/test_conformance/basic/test_image_r8.c +++ /dev/null @@ -1,176 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "harness/compat.h" - -#include -#include -#include -#include - - -#include "procs.h" - -static const char *r_uint8_kernel_code = -"__kernel void test_r_uint8(read_only image2d_t srcimg, __global unsigned char *dst, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" int indx = tid_y * get_image_width(srcimg) + tid_x;\n" -" uint4 color;\n" -"\n" -" color = read_imageui(srcimg, sampler, (int2)(tid_x, tid_y));\n" -" dst[indx] = (unsigned char)(color.x);\n" -"\n" -"}\n"; - - -static unsigned char * -generate_8bit_image(int w, int h, MTdata d) -{ - unsigned char *ptr = (unsigned char*)malloc(w * h * sizeof(unsigned char)); - int i; - - for (i=0; i -#include -#include -#include -#include - - -#include "procs.h" - -static const char *image_dim_kernel_code = -"\n" -"__kernel void test_image_dim(read_only image2d_t srcimg, write_only image2d_t dstimg, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" float4 color;\n" -"\n" -" color = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y));\n" -" write_imagef(dstimg, (int2)(tid_x, tid_y), color);\n" -"\n" -"}\n"; - - -static unsigned char * -generate_8888_image(int w, int h, MTdata d) -{ - unsigned char *ptr = (unsigned char*)malloc(w * h * 4); - int i; - - for (i=0; i (cl_ulong)SIZE_MAX) { - max_mem_size = (cl_ulong)SIZE_MAX; - } - - cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); - test_error(err, "clCreateSampler failed"); - - max_img_width = (int)max_image2d_width; - max_img_height = (int)max_image2d_height; - - // determine max image dim we can allocate - assume RGBA image, 4 bytes per pixel, - // and we want to consume 1/4 of global memory (this is the minimum required to be - // supported by the spec) - max_mem_size /= 4; // use 1/4 - max_mem_size /= 4; // 4 bytes per pixel - max_img_dim = (int)sqrt((double)max_mem_size); - // convert to a power of 2 - { - unsigned int n = (unsigned int)max_img_dim; - unsigned int m = 0x80000000; - - // round-down to the nearest power of 2 - while (m > n) - m >>= 1; - - max_img_dim = (int)m; - } - - if (max_img_width > max_img_dim) - max_img_width = max_img_dim; - if (max_img_height > max_img_dim) - max_img_height = max_img_dim; - - log_info("Adjusted maximum image size to test is %d x %d, which is a max mem size of %gMB.\n", - max_img_width, max_img_height, (max_img_width*max_img_height*4)/(1024.0*1024.0)); - - d = init_genrand( gRandomSeed ); - input_ptr = generate_8888_image(max_img_width, max_img_height, d); - output_ptr = (unsigned char*)malloc(sizeof(unsigned char) * 4 * max_img_width * max_img_height); - - // test power of 2 width, height starting at 1 to 4K - for (i=1,i2=0; i<=max_img_height; i<<=1,i2++) - { - img_height = (1 << i2); - for (j=1,j2=0; j<=max_img_width; j<<=1,j2++) - { - img_width = (1 << j2); - - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_UNORM_INT8; - streams[0] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE), &img_format, img_width, img_height, 0, NULL, NULL); - if (!streams[0]) - { - log_error("create_image_2d failed. width = %d, height = %d\n", img_width, img_height); - free(input_ptr); - free(output_ptr); - free_mtdata(d); - return -1; - } - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_UNORM_INT8; - streams[1] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE), &img_format, img_width, img_height, 0, NULL, NULL); - if (!streams[1]) - { - log_error("create_image_2d failed. width = %d, height = %d\n", img_width, img_height); - clReleaseMemObject(streams[0]); - free(input_ptr); - free(output_ptr); - free_mtdata(d); - return -1; - } - - size_t origin[3] = {0,0,0}; - size_t region[3] = {img_width, img_height, 1}; - err = clEnqueueWriteImage(queue, streams[0], CL_FALSE, origin, region, 0, 0, input_ptr, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clWriteImage failed\n"); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - free(input_ptr); - free(output_ptr); - free_mtdata(d); - return -1; - } - - err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); - err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); - err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler); - if (err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed\n"); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - free(input_ptr); - free(output_ptr); - free_mtdata(d); - return -1; - } - - threads[0] = (size_t)img_width; - threads[1] = (size_t)img_height; - log_info("Testing image dimensions %d x %d with local threads NULL.\n", img_width, img_height); - err = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, NULL, 0, NULL, NULL ); - if (err != CL_SUCCESS) - { - log_error("clEnqueueNDRangeKernel failed\n"); - log_error("Image Dimension test failed. image width = %d, image height = %d, local NULL\n", - img_width, img_height); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - free(input_ptr); - free(output_ptr); - free_mtdata(d); - return -1; - } - err = clEnqueueReadImage(queue, streams[1], CL_TRUE, origin, region, 0, 0, output_ptr, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clReadImage failed\n"); - log_error("Image Dimension test failed. image width = %d, image height = %d, local NULL\n", - img_width, img_height); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - free(input_ptr); - free(output_ptr); - free_mtdata(d); - return -1; - } - err = verify_8888_image(input_ptr, output_ptr, img_width, img_height); - if (err) - { - total_errors++; - log_error("Image Dimension test failed. image width = %d, image height = %d\n", img_width, img_height); - } - - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - } - } - - // cleanup - free(input_ptr); - free(output_ptr); - free_mtdata(d); - clReleaseSampler(sampler); - clReleaseKernel(kernel); - clReleaseProgram(program); - - return total_errors; -} - - - -int -test_imagedim_non_pow2(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems) -{ - cl_mem streams[2]; - cl_image_format img_format; - unsigned char *input_ptr, *output_ptr; - cl_program program; - cl_kernel kernel; - size_t threads[2], local_threads[2]; - cl_ulong max_mem_size; - int img_width, max_img_width; - int img_height, max_img_height; - int max_img_dim; - int i, j, i2, j2, err=0; - size_t max_image2d_width, max_image2d_height; - int total_errors = 0; - size_t max_local_workgroup_size[3]; - MTdata d; - - PASSIVE_REQUIRE_IMAGE_SUPPORT( device ) - - err = create_single_kernel_helper( context, &program, &kernel, 1, &image_dim_kernel_code, "test_image_dim" ); - if (err) - { - log_error("create_program_and_kernel_with_sources failed\n"); - return -1; - } - - size_t work_group_size = 0; - err = clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(work_group_size), &work_group_size, NULL); - test_error(err, "clGetKerenlWorkgroupInfo failed for CL_KERNEL_WORK_GROUP_SIZE"); - - err = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(max_local_workgroup_size), max_local_workgroup_size, NULL); - test_error(err, "clGetDeviceInfo failed for CL_DEVICE_MAX_WORK_ITEM_SIZES"); - - err = clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE,sizeof(max_mem_size), &max_mem_size, NULL); - if (err) - { - log_error("clGetDeviceInfo for CL_DEVICE_GLOBAL_MEM_SIZE failed (%d)\n", err); - return -1; - } - err = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof(max_image2d_width), &max_image2d_width, NULL); - if (err) - { - log_error("clGetDeviceInfo for CL_DEVICE_IMAGE2D_MAX_WIDTH failed (%d)\n", err); - return -1; - } - err = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof(max_image2d_width), &max_image2d_height, NULL); - if (err) - { - log_error("clGetDeviceInfo for CL_DEVICE_IMAGE2D_MAX_HEIGHT failed (%d)\n", err); - return -1; - } - log_info("Device reported max image sizes of %lu x %lu, and max mem size of %gMB.\n", - max_image2d_width, max_image2d_height, max_mem_size/(1024.0*1024.0)); - - cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); - test_error(err, "clCreateSampler failed"); - - max_img_width = (int)max_image2d_width; - max_img_height = (int)max_image2d_height; - - if (max_mem_size > (cl_ulong)SIZE_MAX) { - max_mem_size = (cl_ulong)SIZE_MAX; - } - - // determine max image dim we can allocate - assume RGBA image, 4 bytes per pixel, - // and we want to consume 1/4 of global memory (this is the minimum required to be - // supported by the spec) - max_mem_size /= 4; // use 1/4 - max_mem_size /= 4; // 4 bytes per pixel - max_img_dim = (int)sqrt((double)max_mem_size); - // convert to a power of 2 - { - unsigned int n = (unsigned int)max_img_dim; - unsigned int m = 0x80000000; - - // round-down to the nearest power of 2 - while (m > n) - m >>= 1; - - max_img_dim = (int)m; - } - - if (max_img_width > max_img_dim) - max_img_width = max_img_dim; - if (max_img_height > max_img_dim) - max_img_height = max_img_dim; - - log_info("Adjusted maximum image size to test is %d x %d, which is a max mem size of %gMB.\n", - max_img_width, max_img_height, (max_img_width*max_img_height*4)/(1024.0*1024.0)); - - d = init_genrand( gRandomSeed ); - input_ptr = generate_8888_image(max_img_width, max_img_height, d); - output_ptr = (unsigned char*)malloc(sizeof(unsigned char) * 4 * max_img_width * max_img_height); - - int plus_minus; - for (plus_minus=0; plus_minus < 3; plus_minus++) - { - - // test power of 2 width, height starting at 1 to 4K - for (i=2,i2=1; i<=max_img_height; i<<=1,i2++) - { - img_height = (1 << i2); - for (j=2,j2=1; j<=max_img_width; j<<=1,j2++) - { - img_width = (1 << j2); - - int effective_img_height = img_height; - int effective_img_width = img_width; - - local_threads[0] = 1; - local_threads[1] = 1; - - switch (plus_minus) { - case 0: - effective_img_height--; - local_threads[0] = work_group_size > max_local_workgroup_size[0] ? max_local_workgroup_size[0] : work_group_size; - while (img_width%local_threads[0] != 0) - local_threads[0]--; - break; - case 1: - effective_img_width--; - local_threads[1] = work_group_size > max_local_workgroup_size[1] ? max_local_workgroup_size[1] : work_group_size; - while (img_height%local_threads[1] != 0) - local_threads[1]--; - break; - case 2: - effective_img_width--; - effective_img_height--; - break; - default: - break; - } - - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_UNORM_INT8; - streams[0] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE), &img_format, effective_img_width, effective_img_height, 0, NULL, NULL); - if (!streams[0]) - { - log_error("create_image_2d failed. width = %d, height = %d\n", effective_img_width, effective_img_height); - free(input_ptr); - free(output_ptr); - free_mtdata(d); - return -1; - } - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_UNORM_INT8; - streams[1] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE), &img_format, effective_img_width, effective_img_height, 0, NULL, NULL); - if (!streams[1]) - { - log_error("create_image_2d failed. width = %d, height = %d\n", effective_img_width, effective_img_height); - clReleaseMemObject(streams[0]); - free(input_ptr); - free(output_ptr); - free_mtdata(d); - return -1; - } - - size_t origin[3] = {0,0,0}; - size_t region[3] = {effective_img_width, effective_img_height, 1}; - err = clEnqueueWriteImage(queue, streams[0], CL_FALSE, origin, region, 0, 0, input_ptr, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clWriteImage failed\n"); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - free(input_ptr); - free(output_ptr); - free_mtdata(d); - return -1; - } - - err = clSetKernelArg(kernel, 0, sizeof streams[0], &streams[0]); - err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]); - err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler); - if (err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed\n"); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - free(input_ptr); - free(output_ptr); - free_mtdata(d); - return -1; - } - - threads[0] = (size_t)effective_img_width; - threads[1] = (size_t)effective_img_height; - log_info("Testing image dimensions %d x %d with local threads %d x %d.\n", - effective_img_width, effective_img_height, (int)local_threads[0], (int)local_threads[1]); - err = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, local_threads, 0, NULL, NULL ); - if (err != CL_SUCCESS) - { - log_error("clEnqueueNDRangeKernel failed\n"); - log_error("Image Dimension test failed. image width = %d, image height = %d, local %d x %d\n", - effective_img_width, effective_img_height, (int)local_threads[0], (int)local_threads[1]); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - free(input_ptr); - free(output_ptr); - free_mtdata(d); - return -1; - } - err = clEnqueueReadImage(queue, streams[1], CL_TRUE, origin, region, 0, 0, output_ptr, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clReadImage failed\n"); - log_error("Image Dimension test failed. image width = %d, image height = %d, local %d x %d\n", - effective_img_width, effective_img_height, (int)local_threads[0], (int)local_threads[1]); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - free(input_ptr); - free(output_ptr); - free_mtdata(d); - return -1; - } - err = verify_8888_image(input_ptr, output_ptr, effective_img_width, effective_img_height); - if (err) - { - total_errors++; - log_error("Image Dimension test failed. image width = %d, image height = %d\n", effective_img_width, effective_img_height); - } - - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - } - } - - } - - // cleanup - free(input_ptr); - free(output_ptr); - free_mtdata(d); - clReleaseSampler(sampler); - clReleaseKernel(kernel); - clReleaseProgram(program); - - return total_errors; -} - - - - diff --git a/test_conformance/compatibility/test_conformance/basic/test_imagenpot.c b/test_conformance/compatibility/test_conformance/basic/test_imagenpot.c deleted file mode 100644 index 39126afd..00000000 --- a/test_conformance/compatibility/test_conformance/basic/test_imagenpot.c +++ /dev/null @@ -1,220 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "harness/compat.h" - -#include -#include -#include -#include - - -#include "procs.h" - -static const char *rgba8888_kernel_code = -"\n" -"__kernel void test_rgba8888(read_only image2d_t srcimg, write_only image2d_t dstimg, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" float4 color;\n" -"\n" -" if ( (tid_x >= get_image_width(dstimg)) || (tid_y >= get_image_height(dstimg)) )\n" -" return;\n" -" color = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y));\n" -" write_imagef(dstimg, (int2)(tid_x, tid_y), color);\n" -"\n" -"}\n"; - - -static unsigned char * -generate_8888_image(int w, int h, MTdata d) -{ - unsigned char *ptr = (unsigned char*)malloc(w * h * 4); - int i; - - for (i=0; i max_local_workgroup_size[0]) - local_workgroup_size = max_local_workgroup_size[0]; - - global_threads[0] = ((img_width + local_workgroup_size - 1) / local_workgroup_size) * local_workgroup_size; - global_threads[1] = img_height; - local_threads[0] = local_workgroup_size; - local_threads[1] = 1; - err = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, global_threads, local_threads, 0, NULL, NULL ); - - if (err != CL_SUCCESS) - { - log_error("%s clEnqueueNDRangeKernel failed\n", __FUNCTION__); - free_mtdata(d); - return -1; - } - err = clEnqueueReadImage(queue, streams[1], CL_TRUE, - origin, region, 0, 0, - (void *)output_ptr, - 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clEnqueueReadBuffer failed\n"); - return -1; - } - - err = verify_rgba8888_image(input_ptr, output_ptr, img_width, img_height); - - // cleanup - clReleaseSampler(sampler); - clReleaseMemObject(streams[0]); - clReleaseMemObject(streams[1]); - clReleaseKernel(kernel); - clReleaseProgram(program); - free(input_ptr); - free(output_ptr); - - if (err) - break; - } - - free_mtdata(d); - - return err; -} - - - - - diff --git a/test_conformance/compatibility/test_conformance/basic/test_multireadimagemultifmt.c b/test_conformance/compatibility/test_conformance/basic/test_multireadimagemultifmt.c deleted file mode 100644 index 895aeb55..00000000 --- a/test_conformance/compatibility/test_conformance/basic/test_multireadimagemultifmt.c +++ /dev/null @@ -1,230 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "harness/compat.h" - -#include -#include -#include -#include - - -#include "procs.h" - -static const char *multireadimage_kernel_code = -"__kernel void test_multireadimage(read_only image2d_t img0, read_only image2d_t img1, \n" -" read_only image2d_t img2, __global float4 *dst, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" int2 tid = (int2)(tid_x, tid_y);\n" -" int indx = tid_y * get_image_width(img1) + tid_x;\n" -" float4 sum;\n" -"\n" -" sum = read_imagef(img0, sampler, tid);\n" -" sum += read_imagef(img1, sampler, tid);\n" -" sum += read_imagef(img2, sampler, tid);\n" -"\n" -" dst[indx] = sum;\n" -"}\n"; - -#define MAX_ERR 1e-7f - -static unsigned char * -generate_8888_image(int w, int h, MTdata d) -{ - unsigned char *ptr = (unsigned char*)malloc(w * h * 4); - int i; - - for (i=0; i max_ulp) - max_ulp = ulp; - } - - if (max_ulp > max_ulp_allowed) { - log_error("READ_MULTIREADIMAGE_MULTIFORMAT test failed. Max ulp error = %g\n", max_ulp); - return -1; - } - - log_info("READ_MULTIREADIMAGE_MULTIFORMAT test passed. Max ulp error = %g\n", max_ulp); - return 0; -} - - -int -test_mri_multiple(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) -{ - cl_mem streams[4]; - cl_image_format img_format; - void *input_ptr[3], *output_ptr; - cl_program program; - cl_kernel kernel; - size_t threads[2]; - int img_width = 512; - int img_height = 512; - int i, err; - MTdata d; - - PASSIVE_REQUIRE_IMAGE_SUPPORT( device ) - - d = init_genrand( gRandomSeed ); - input_ptr[0] = (void *)generate_8888_image(img_width, img_height, d); - input_ptr[1] = (void *)generate_16bit_image(img_width, img_height, d); - input_ptr[2] = (void *)generate_float_image(img_width, img_height, d); - free_mtdata(d); d = NULL; - - output_ptr = (void *)malloc(sizeof(float) * 4 * img_width * img_height); - - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_UNORM_INT8; - streams[0] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE), &img_format, img_width, img_height, 0, NULL, NULL); - if (!streams[0]) - { - log_error("create_image_2d failed\n"); - return -1; - } - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_UNORM_INT16; - streams[1] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE), &img_format, img_width, img_height, 0, NULL, NULL); - if (!streams[1]) - { - log_error("create_image_2d failed\n"); - return -1; - } - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_FLOAT; - streams[2] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE), &img_format, img_width, img_height, 0, NULL, NULL); - if (!streams[2]) - { - log_error("create_image_2d failed\n"); - return -1; - } - - streams[3] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(float)*4 * img_width*img_height, NULL, NULL); - if (!streams[3]) - { - log_error("clCreateBuffer failed\n"); - return -1; - } - - for (i=0; i<3; i++) - { - size_t origin[3] = {0,0,0}, region[3]={img_width, img_height,1}; - err = clEnqueueWriteImage(queue, streams[i], CL_TRUE, origin, region, 0, 0, input_ptr[i], 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clWriteImage failed\n"); - return -1; - } - } - - err = create_single_kernel_helper( context, &program, &kernel, 1, &multireadimage_kernel_code, "test_multireadimage"); - if (err) - return -1; - - cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); - test_error(err, "clCreateSampler failed"); - - for (i=0; i<4; i++) - err |= clSetKernelArg(kernel, i,sizeof streams[i], &streams[i]); - err |= clSetKernelArg(kernel, 4, sizeof sampler, &sampler); - - if (err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed\n"); - return -1; - } - - threads[0] = (size_t)img_width; - threads[1] = (size_t)img_height; - - err = clEnqueueNDRangeKernel( queue, kernel, 2, NULL, threads, NULL, 0, NULL, NULL ); - if (err != CL_SUCCESS) - { - log_error("clEnqueueNDRangeKernel failed\n"); - return -1; - } - err = clEnqueueReadBuffer( queue, streams[3], CL_TRUE, 0, sizeof(float)*4*img_width*img_height, (void *)output_ptr, 0, NULL, NULL ); - if (err != CL_SUCCESS) - { - log_error("clEnqueueReadBuffer failed\n"); - return -1; - } - - err = verify_multireadimage(input_ptr, (float*)output_ptr, img_width, img_height); - - // cleanup - clReleaseSampler(sampler); - for (i=0; i<4; i++) - clReleaseMemObject(streams[i]); - clReleaseKernel(kernel); - clReleaseProgram(program); - for (i=0; i<3; i++) - free(input_ptr[i]); - free(output_ptr); - - return err; -} - - - - - diff --git a/test_conformance/compatibility/test_conformance/basic/test_multireadimageonefmt.c b/test_conformance/compatibility/test_conformance/basic/test_multireadimageonefmt.c deleted file mode 100644 index 91c2b5a8..00000000 --- a/test_conformance/compatibility/test_conformance/basic/test_multireadimageonefmt.c +++ /dev/null @@ -1,198 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "harness/compat.h" - -#include -#include -#include -#include - - -#include "procs.h" - -static const char *multireadimage_kernel_code = -"__kernel void test_multireadimage(int n, int m, sampler_t sampler, \n" -" read_only image2d_t img0, read_only image2d_t img1, \n" -" read_only image2d_t img2, read_only image2d_t img3, \n" -" read_only image2d_t img4, read_only image2d_t img5, \n" -" read_only image2d_t img6, __global float4 *dst)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" int2 tid = (int2)(tid_x, tid_y);\n" -" int indx = tid_y * get_image_width(img5) + tid_x;\n" -" float4 sum;\n" -"\n" -" sum = read_imagef(img0, sampler, tid);\n" -" sum += read_imagef(img1, sampler, tid);\n" -" sum += read_imagef(img2, sampler, tid);\n" -" sum += read_imagef(img3, sampler, tid);\n" -" sum += read_imagef(img4, sampler, tid);\n" -" sum += read_imagef(img5, sampler, tid);\n" -" sum += read_imagef(img6, sampler, tid);\n" -"\n" -" dst[indx] = sum;\n" -"}\n"; - - -static unsigned char * -generate_8888_image(int w, int h, MTdata d) -{ - unsigned char *ptr = (unsigned char*)malloc(w * h * 4); - int i; - - for (i=0; i max_ulp) - max_ulp = ulp; - } - - if (max_ulp > max_ulp_allowed) - { - log_error("READ_MULTIREADIMAGE_RGBA8888 test failed. Max ULP err = %g\n", max_ulp); - return -1; - } - log_info("READ_MULTIREADIMAGE_RGBA8888 test passed. Max ULP err = %g\n", max_ulp); - return 0; -} - - -int test_mri_one(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) -{ - cl_mem streams[8]; - cl_image_format img_format; - void *input_ptr[7], *output_ptr; - cl_program program; - cl_kernel kernel; - size_t threads[2]; - int img_width = 512; - int img_height = 512; - int i, err; - size_t origin[3] = {0, 0, 0}; - size_t region[3] = {img_width, img_height, 1}; - size_t length = img_width * img_height * 4 * sizeof(float); - MTdata d; - - PASSIVE_REQUIRE_IMAGE_SUPPORT( device ) - - output_ptr = malloc(length); - - d = init_genrand( gRandomSeed ); - for (i=0; i<7; i++) { - input_ptr[i] = (void *)generate_8888_image(img_width, img_height, d); - - img_format.image_channel_order = CL_RGBA; - img_format.image_channel_data_type = CL_UNORM_INT8; - streams[i] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format, img_width, img_height, 0, NULL, NULL); - if (!streams[i]) - { - log_error("create_image_2d failed\n"); - return -1; - } - - err = clEnqueueWriteImage(queue, streams[i], CL_TRUE, origin, region, 0, 0, input_ptr[i], 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clWriteImage failed\n"); - return -1; - } - } - free_mtdata(d); d = NULL; - - - streams[7] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL); - if (!streams[7]) - { - log_error("clCreateArray failed\n"); - return -1; - } - - err = create_single_kernel_helper(context, &program, &kernel, 1, &multireadimage_kernel_code, "test_multireadimage"); - if (err) - return -1; - - cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err); - test_error(err, "clCreateSampler failed"); - - err = clSetKernelArg(kernel, 0, sizeof i, &i); - err |= clSetKernelArg(kernel, 1, sizeof err, &err); - err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler); - for (i=0; i<8; i++) - err |= clSetKernelArg(kernel, 3+i, sizeof streams[i], &streams[i]); - - if (err != CL_SUCCESS) - { - log_error("clSetKernelArgs failed\n"); - return -1; - } - - threads[0] = (unsigned int)img_width; - threads[1] = (unsigned int)img_height; - - err = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, threads, NULL, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clExecuteKernel failed\n"); - return -1; - } - err = clEnqueueReadBuffer(queue, streams[7], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - log_error("clReadArray failed\n"); - return -1; - } - - err = verify_multireadimage(input_ptr, 7, (float *)output_ptr, img_width, img_height); - - // cleanup - clReleaseSampler(sampler); - for (i=0; i<8; i++) - clReleaseMemObject(streams[i]); - clReleaseKernel(kernel); - clReleaseProgram(program); - for (i=0; i<7; i++) - free(input_ptr[i]); - free(output_ptr); - - return err; -} - - - - - diff --git a/test_conformance/compatibility/test_conformance/basic/test_readimage3d.c b/test_conformance/compatibility/test_conformance/basic/test_readimage3d.c deleted file mode 100644 index ad6f29e2..00000000 --- a/test_conformance/compatibility/test_conformance/basic/test_readimage3d.c +++ /dev/null @@ -1,230 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "harness/compat.h" - -#include -#include -#include -#include - - -#include "procs.h" - -static const char *bgra8888_kernel_code = -"\n" -"__kernel void test_bgra8888(read_only image3d_t srcimg, __global float4 *dst, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" int tid_z = get_global_id(2);\n" -" int indx = (tid_z * get_image_height(srcimg) + tid_y) * get_image_width(srcimg) + tid_x;\n" -" float4 color;\n" -"\n" -" color = read_imagef(srcimg, sampler, (int4)(tid_x, tid_y, tid_z, 0));\n" -" dst[indx].x = color.z;\n" -" dst[indx].y = color.y;\n" -" dst[indx].z = color.x;\n" -" dst[indx].w = color.w;\n" -"\n" -"}\n"; - - -static const char *rgba8888_kernel_code = -"\n" -"__kernel void test_rgba8888(read_only image3d_t srcimg, __global float4 *dst, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" int tid_z = get_global_id(2);\n" -" int indx = (tid_z * get_image_height(srcimg) + tid_y) * get_image_width(srcimg) + tid_x;\n" -" float4 color;\n" -"\n" -" color = read_imagef(srcimg, sampler, (int4)(tid_x, tid_y, tid_z, 0));\n" -" //indx *= 4;\n" -" dst[indx].x = color.x;\n" -" dst[indx].y = color.y;\n" -" dst[indx].z = color.z;\n" -" dst[indx].w = color.w;\n" -"\n" -"}\n"; - - -static unsigned char * -generate_3d_image8(int w, int h, int d, MTdata data) -{ - unsigned char *ptr = (unsigned char*)malloc(w * h * d * 4); - int i; - - for (i=0; i -#include -#include -#include - - -#include "procs.h" - - -static const char *rgbaFFFF_kernel_code = -"__kernel void test_rgbaFFFF(read_only image3d_t srcimg, __global float *dst, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" int tid_z = get_global_id(2);\n" -" int indx = (tid_z * get_image_height(srcimg) + tid_y) * get_image_width(srcimg) + tid_x;\n" -" float4 color;\n" -"\n" -" color = read_imagef(srcimg, sampler, (int4)(tid_x, tid_y, tid_z, 0));\n" -" indx *= 4;\n" -" dst[indx+0] = color.x;\n" -" dst[indx+1] = color.y;\n" -" dst[indx+2] = color.z;\n" -" dst[indx+3] = color.w;\n" -"\n" -"}\n"; - - -static float * -generate_float_image(int w, int h, int d, MTdata data) -{ - float *ptr = (float*)malloc(w * h * d * 4 * sizeof(float)); - int i; - - for (i=0; i -#include -#include -#include - - -#include "procs.h" - -static const char *rgba16_kernel_code = -"__kernel void test_rgba16(read_only image3d_t srcimg, __global ushort4 *dst, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" int tid_z = get_global_id(2);\n" -" int indx = (tid_z * get_image_height(srcimg) + tid_y) * get_image_width(srcimg) + tid_x;\n" -" float4 color;\n" -"\n" -" color = read_imagef(srcimg, sampler, (int4)(tid_x, tid_y, tid_z, 0));\n" -" ushort4 dst_write;\n" -" dst_write.x = convert_ushort_rte(color.x * 65535.0f);\n" -" dst_write.y = convert_ushort_rte(color.y * 65535.0f);\n" -" dst_write.z = convert_ushort_rte(color.z * 65535.0f);\n" -" dst_write.w = convert_ushort_rte(color.w * 65535.0f);\n" -" dst[indx] = dst_write;\n" -"\n" -"}\n"; - - -static unsigned short * -generate_16bit_image(int w, int h, int d, MTdata data) -{ - unsigned short *ptr = (cl_ushort*)malloc(w * h * d * 4 * sizeof(cl_ushort)); - int i; - - for (i=0; i -#include -#include -#include - - -#include "procs.h" - - -static const char *rgbaFFFF_kernel_code = -"__kernel void test_rgbaFFFF(read_only image2d_t srcimg, __global float *dst, sampler_t smp)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" int indx = tid_y * get_image_width(srcimg) + tid_x;\n" -" float4 color;\n" -"\n" -" color = read_imagef(srcimg, smp, (int2)(tid_x, tid_y));\n" -" indx *= 4;\n" -" dst[indx+0] = color.x;\n" -" dst[indx+1] = color.y;\n" -" dst[indx+2] = color.z;\n" -" dst[indx+3] = color.w;\n" -"\n" -"}\n"; - - -static float * -generate_float_image(int w, int h, MTdata d) -{ - float *ptr = (float*)malloc(w * h * 4 * sizeof(float)); - int i; - - for (i=0; i -#include -#include -#include - - -#include "procs.h" - -static const char *rgba16_kernel_code = -"__kernel void test_rgba16(read_only image2d_t srcimg, __global ushort4 *dst, sampler_t smp)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" int indx = tid_y * get_image_width(srcimg) + tid_x;\n" -" float4 color;\n" -"\n" -" color = read_imagef(srcimg, smp, (int2)(tid_x, tid_y));\n" -" ushort4 dst_write;\n" -" dst_write.x = convert_ushort_rte(color.x * 65535.0f);\n" -" dst_write.y = convert_ushort_rte(color.y * 65535.0f);\n" -" dst_write.z = convert_ushort_rte(color.z * 65535.0f);\n" -" dst_write.w = convert_ushort_rte(color.w * 65535.0f);\n" -" dst[indx] = dst_write;\n" -"\n" -"}\n"; - - -static unsigned short * -generate_16bit_image(int w, int h, MTdata d) -{ - cl_ushort *ptr = (cl_ushort*)malloc(w * h * 4 * sizeof(cl_ushort)); - int i; - - for (i=0; i