ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-20 06:29:02 +00:00
Code Packages Projects Releases Wiki Activity

Use old-style sampler creation in basic suite and remove duplicate compatibility tests (#510)

Browse Source 
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 <kpet@free.fr>
This commit is contained in:
Kévin Petit
2019-12-11 10:36:45 +00:00
committed by GitHub
parent 7127b2a05f
commit 1c04050b5b
26 changed files with 32 additions and 3212 deletions
Download Patch File Download Diff File Expand all files Collapse all files

18
test_conformance/basic/test_image_multipass.c
View File

@@ -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};

9
test_conformance/basic/test_image_param.c
View File

@@ -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++ )

9
test_conformance/basic/test_image_r8.c
View File

@@ -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]);

24
test_conformance/basic/test_imagedim.c
View File

@@ -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;

9
test_conformance/basic/test_imagenpot.c
View File

@@ -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]);

9
test_conformance/basic/test_multireadimagemultifmt.c
View File

@@ -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]);

9
test_conformance/basic/test_multireadimageonefmt.c
View File

@@ -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);

9
test_conformance/basic/test_readimage3d.c
View File

@@ -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]);

9
test_conformance/basic/test_readimage3d_fp32.c
View File

@@ -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]);

9
test_conformance/basic/test_readimage3d_int16.c
View File

@@ -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]);

9
test_conformance/basic/test_readimage_fp32.c
View File

@@ -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]);

9
test_conformance/basic/test_readimage_int16.c
View File

@@ -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]);

9
test_conformance/compatibility/test_conformance/basic/CMakeLists.txt
View File

@@ -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

15
test_conformance/compatibility/test_conformance/basic/main.c
View File

@@ -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 ),

643
test_conformance/compatibility/test_conformance/basic/test_image_multipass.c
View File

@@ -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 <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#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;
}

251
test_conformance/compatibility/test_conformance/basic/test_image_param.c
View File

@@ -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 <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#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<char> 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<cl_float> 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;
}

176
test_conformance/compatibility/test_conformance/basic/test_image_r8.c
View File

@@ -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 <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#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<w*h; i++)
ptr[i] = (unsigned char)genrand_int32(d);
return ptr;
}
static int
verify_8bit_image(unsigned char *image, unsigned char *outptr, int w, int h)
{
int i;
for (i=0; i<w*h; i++)
{
if (outptr[i] != image[i])
{
log_error("READ_IMAGE_R_UNSIGNED_INT8 test failed\n");
return -1;
}
}
log_info("READ_IMAGE_R_UNSIGNED_INT8 test passed\n");
return 0;
}
int
test_image_r8(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
cl_mem streams[2];
cl_image_format img_format;
cl_uchar *input_ptr, *output_ptr;
cl_program program;
cl_kernel kernel;
size_t threads[3];
int img_width = 512;
int img_height = 512;
int err;
MTdata d;
PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
img_format.image_channel_order = CL_R;
img_format.image_channel_data_type = CL_UNSIGNED_INT8;
// early out if this image type is not supported
if( ! is_image_format_supported( context, (cl_mem_flags)(CL_MEM_READ_ONLY), CL_MEM_OBJECT_IMAGE2D, &img_format ) ) {
log_info("WARNING: Image type not supported; skipping test.\n");
return 0;
}
d = init_genrand( gRandomSeed );
input_ptr = generate_8bit_image(img_width, img_height, d);
free_mtdata(d); d = NULL;
output_ptr = (cl_uchar*)malloc(sizeof(cl_uchar) * img_width * img_height);
streams[0] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_ONLY), &img_format, img_width, img_height, 0, NULL, NULL);
if (!streams[0])
{
log_error("create_image_2d failed\n");
return -1;
}
streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uchar) * img_width*img_height, NULL, NULL);
if (!streams[1])
{
log_error("clCreateBuffer failed\n");
return -1;
}
size_t origin[3] = {0,0,0}, region[3]={img_width, img_height, 1};
err = clEnqueueWriteImage(queue, streams[0], CL_TRUE,
origin, region, 0, 0,
input_ptr,
0, NULL, NULL);
if (err != CL_SUCCESS)
{
log_error("clWriteImage failed: %d\n", err);
return -1;
}
err = create_single_kernel_helper(context, &program, &kernel, 1, &r_uint8_kernel_code, "test_r_uint8" );
if (err) {
log_error("Failed to create kernel and program: %d\n", 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 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: %d\n", err);
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[1], CL_TRUE, 0, sizeof(cl_uchar)*img_width*img_height, (void *)output_ptr, 0, NULL, NULL );
if (err != CL_SUCCESS)
{
log_error("clEnqueueReadBuffer failed\n");
return -1;
}
err = verify_8bit_image(input_ptr, output_ptr, img_width, img_height);
// cleanup
clReleaseMemObject(streams[0]);
clReleaseMemObject(streams[1]);
clReleaseKernel(kernel);
clReleaseProgram(program);
free(input_ptr);
free(output_ptr);
return err;
}

514
test_conformance/compatibility/test_conformance/basic/test_imagedim.c
View File

@@ -1,514 +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 <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#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<w*h*4; i++)
ptr[i] = (unsigned char)genrand_int32(d);
return ptr;
}
static int
verify_8888_image(unsigned char *image, unsigned char *outptr, int w, int h)
{
int i;
for (i=0; i<w*h; i++)
{
if (outptr[i] != image[i])
return -1;
}
return 0;
}
int
test_imagedim_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];
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;
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;
}
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));
if (max_mem_size > (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;
}

220
test_conformance/compatibility/test_conformance/basic/test_imagenpot.c
View File

@@ -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 <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#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<w*h*4; i++)
ptr[i] = (unsigned char)genrand_int32(d);
return ptr;
}
static int
verify_rgba8888_image(unsigned char *src, unsigned char *dst, int w, int h)
{
int i;
for (i=0; i<w*h*4; i++)
{
if (dst[i] != src[i])
{
log_error("NPOT_IMAGE_RGBA_UNORM_INT8 test for width = %d, height = %d failed\n", w, h);
return -1;
}
}
log_info("NPOT_IMAGE_RGBA_UNORM_INT8 test for width = %d, height = %d passed\n", w, h);
return 0;
}
int img_width_selection[] = { 97, 111, 322, 479 };
int img_height_selection[] = { 149, 222, 754, 385 };
int
test_imagenpot(cl_device_id device_id, cl_context context, cl_command_queue queue, int num_elements)
{
cl_mem streams[2];
cl_image_format img_format;
unsigned char *input_ptr, *output_ptr;
cl_program program;
cl_kernel kernel;
size_t global_threads[3], local_threads[3];
size_t local_workgroup_size;
int img_width;
int img_height;
int err;
cl_uint m;
size_t max_local_workgroup_size[3];
MTdata d;
PASSIVE_REQUIRE_IMAGE_SUPPORT( device_id )
cl_device_type device_type;
err = clGetDeviceInfo(device_id, CL_DEVICE_TYPE, sizeof(device_type), &device_type, NULL);
if (err) {
log_error("Failed to get device type: %d\n",err);
return -1;
}
d = init_genrand( gRandomSeed );
for (m=0; m<sizeof(img_width_selection)/sizeof(int); m++)
{
img_width = img_width_selection[m];
img_height = img_height_selection[m];
input_ptr = generate_8888_image(img_width, img_height, d);
output_ptr = (unsigned char*)malloc(sizeof(unsigned char) * 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");
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\n");
free_mtdata(d);
return -1;
}
size_t origin[3] = {0,0,0}, region[3] = {img_width, img_height, 1};
err = clEnqueueWriteImage(queue, streams[0], CL_TRUE,
origin, region, 0, 0,
input_ptr,
0, NULL, NULL);
if (err != CL_SUCCESS)
{
log_error("clWriteImage failed\n");
free_mtdata(d);
return -1;
}
err = create_single_kernel_helper(context, &program, &kernel, 1, &rgba8888_kernel_code, "test_rgba8888" );
if (err)
{
log_error("Failed to create kernel and program: %d\n", err);
free_mtdata(d);
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 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");
free_mtdata(d);
return -1;
}
err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local_workgroup_size), &local_workgroup_size, NULL);
test_error(err, "clGetKernelWorkGroupInfo for CL_KERNEL_WORK_GROUP_SIZE failed");
err = clGetDeviceInfo(device_id, 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");
// Pick the minimum of the device and the kernel
if (local_workgroup_size > 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;
}

230
test_conformance/compatibility/test_conformance/basic/test_multireadimagemultifmt.c
View File

@@ -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 <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#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<w*h*4; i++)
ptr[i] = (unsigned char)genrand_int32(d);
return ptr;
}
static unsigned short *
generate_16bit_image(int w, int h, MTdata d)
{
unsigned short *ptr = (unsigned short*)malloc(w * h * 4 * sizeof(unsigned short));
int i;
for (i=0; i<w*h*4; i++)
ptr[i] = (unsigned short)genrand_int32(d);
return ptr;
}
static float *
generate_float_image(int w, int h, MTdata d)
{
float *ptr = (float*)malloc(w * h * 4 * (int)sizeof(float));
int i;
for (i=0; i<w*h*4; i++)
ptr[i] = get_random_float(-0x40000000, 0x40000000, d);
return ptr;
}
static int
verify_multireadimage(void *image[], float *outptr, int w, int h)
{
int i;
float sum;
float ulp, max_ulp = 0.0f;
// ULP error of 1.5 for each read_imagef plus 0.5 for each addition.
float max_ulp_allowed = (float)(3*1.5+2*0.5);
for (i=0; i<w*h*4; i++)
{
sum = (float)((unsigned char *)image[0])[i] / 255.0f;
sum += (float)((unsigned short *)image[1])[i] / 65535.0f;
sum += (float)((float *)image[2])[i];
ulp = Ulp_Error(outptr[i], sum);
if (ulp > 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;
}

198
test_conformance/compatibility/test_conformance/basic/test_multireadimageonefmt.c
View File

@@ -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 <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#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<w*h*4; i++)
ptr[i] = (unsigned char)genrand_int32(d);
return ptr;
}
static int
verify_multireadimage(void *image[], int num_images, float *outptr, int w, int h)
{
int i, j;
float sum;
float ulp, max_ulp = 0.0f;
// ULP error of 1.5 for each read_imagef plus 0.5 for each addition.
float max_ulp_allowed = (float)(num_images*1.5+0.5*(num_images-1));
for (i=0; i<w*h*4; i++)
{
sum = 0.0f;
for (j=0; j<num_images; j++)
{
sum += ((float)((unsigned char *)image[j])[i] / 255.0f);
}
ulp = Ulp_Error(outptr[i], sum);
if (ulp > 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;
}

230
test_conformance/compatibility/test_conformance/basic/test_readimage3d.c
View File

@@ -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 <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#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<w*h*d*4; i++)
ptr[i] = (unsigned char)genrand_int32(data);
return ptr;
}
static int
verify_3d_image8(double *image, float *outptr, int w, int h, int d)
{
int i;
for (i=0; i<w*h*d*4; i++)
{
if (outptr[i] != (float)image[i])
{
float ulps = Ulp_Error( outptr[i], image[i]);
if(! (fabsf(ulps) < 1.5f) )
{
log_error( "ERROR: Data sample %d does not validate! Expected (%a), got (%a), ulp %f\n",
(int)i, image[i], outptr[ i ], ulps );
return -1;
}
}
}
return 0;
}
static double *
prepare_reference(unsigned char * input_ptr, int w, int h, int d)
{
double *ptr = (double*)malloc(w * h * d * 4 * sizeof(double));
int i;
for (i=0; i<w*h*d*4; i++)
ptr[i] = ((double)input_ptr[i]/255);
return ptr;
}
int test_readimage3d(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
cl_mem streams[3];
cl_program program[2];
cl_kernel kernel[2];
cl_image_format img_format;
unsigned char *input_ptr[2];
float *output_ptr;
double *ref_ptr[2];
size_t threads[3];
int img_width = 64;
int img_height = 64;
int img_depth = 64;
int i, err;
size_t origin[3] = {0, 0, 0};
size_t region[3] = {img_width, img_height, img_depth};
size_t length = img_width * img_height * img_depth * 4 * sizeof(float);
PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( device )
MTdata d = init_genrand( gRandomSeed );
input_ptr[0] = generate_3d_image8(img_width, img_height, img_depth, d);
input_ptr[1] = generate_3d_image8(img_width, img_height, img_depth, d);
ref_ptr[0] = prepare_reference(input_ptr[0], img_width, img_height, img_depth);
ref_ptr[1] = prepare_reference(input_ptr[1], img_width, img_height, img_depth);
free_mtdata(d); d = NULL;
output_ptr = (float*)malloc(length);
img_format.image_channel_order = CL_BGRA;
img_format.image_channel_data_type = CL_UNORM_INT8;
streams[0] = create_image_3d(context, CL_MEM_READ_ONLY, &img_format, img_width, img_height, img_depth, 0, 0, NULL, &err);
test_error(err, "create_image_3d failed");
img_format.image_channel_order = CL_RGBA;
img_format.image_channel_data_type = CL_UNORM_INT8;
streams[1] = create_image_3d(context, CL_MEM_READ_ONLY, &img_format, img_width, img_height, img_depth, 0, 0, NULL, &err);
test_error(err, "create_image_3d failed");
streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
test_error(err, "clCreateBuffer failed");
err = clEnqueueWriteImage(queue, streams[0], CL_TRUE, origin, region, 0, 0, input_ptr[0], 0, NULL, NULL);
test_error(err, "clEnqueueWriteImage failed");
err = clEnqueueWriteImage(queue, streams[1], CL_TRUE, origin, region, 0, 0, input_ptr[1], 0, NULL, NULL);
test_error(err, "clEnqueueWriteImage failed");
err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &bgra8888_kernel_code, "test_bgra8888" );
if (err)
return -1;
err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &rgba8888_kernel_code, "test_rgba8888" );
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], 0, sizeof streams[0], &streams[0]);
err |= clSetKernelArg(kernel[0], 1, sizeof streams[2], &streams[2]);
err |= clSetKernelArg(kernel[0], 2, sizeof sampler, &sampler);
test_error(err, "clSetKernelArg failed");
err = clSetKernelArg(kernel[1], 0, sizeof streams[1], &streams[1]);
err |= clSetKernelArg(kernel[1], 1, sizeof streams[2], &streams[2]);
err |= clSetKernelArg(kernel[1], 2, sizeof sampler, &sampler);
test_error(err, "clSetKernelArg failed");
threads[0] = (unsigned int)img_width;
threads[1] = (unsigned int)img_height;
threads[2] = (unsigned int)img_depth;
for (i=0; i<2; i++)
{
err = clEnqueueNDRangeKernel(queue, kernel[i], 3, NULL, threads, NULL, 0, NULL, NULL);
test_error(err, "clEnqueueNDRangeKernel failed");
err = clEnqueueReadBuffer(queue, streams[2], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
test_error(err, "clEnqueueReadBuffer failed");
switch (i)
{
case 0:
err = verify_3d_image8(ref_ptr[i], output_ptr, img_width, img_height, img_depth);
if ( err != 0 )
log_info("READ_IMAGE3D_BGRA_UNORM_INT8 test passed\n");
break;
case 1:
err = verify_3d_image8(ref_ptr[i], output_ptr, img_width, img_height, img_depth);
if ( err != 0 )
log_info("READ_IMAGE3D_RGBA_UNORM_INT8 test passed\n");
break;
}
if (err)
break;
}
// cleanup
clReleaseSampler(sampler);
clReleaseMemObject(streams[0]);
clReleaseMemObject(streams[1]);
clReleaseMemObject(streams[2]);
for (i=0; i<2; i++)
{
clReleaseKernel(kernel[i]);
clReleaseProgram(program[i]);
}
free(input_ptr[0]);
free(input_ptr[1]);
free(output_ptr);
free(ref_ptr[0]);
free(ref_ptr[1]);
return err;
}

147
test_conformance/compatibility/test_conformance/basic/test_readimage3d_fp32.c
View File

@@ -1,147 +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 <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#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<w*h*d*4; i++)
ptr[i] = get_random_float(-0x40000000, 0x40000000, data);
return ptr;
}
static int
verify_float_image(float *image, float *outptr, int w, int h, int d)
{
int i;
for (i=0; i<w*h*d*4; i++)
{
if (outptr[i] != image[i])
{
log_error("READ_IMAGE3D_RGBA_FLOAT test failed\n");
return -1;
}
}
log_info("READ_IMAGE3D_RGBA_FLOAT test passed\n");
return 0;
}
int test_readimage3d_fp32(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
cl_mem streams[2];
cl_program program;
cl_kernel kernel;
cl_image_format img_format;
float *input_ptr, *output_ptr;
size_t threads[3];
int img_width = 64;
int img_height = 64;
int img_depth = 64;
int err;
size_t origin[3] = {0, 0, 0};
size_t region[3] = {img_width, img_height, img_depth};
size_t length = img_width * img_height * img_depth * 4 * sizeof(float);
PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( device )
MTdata d = init_genrand( gRandomSeed );
input_ptr = generate_float_image(img_width, img_height, img_depth, d);
free_mtdata(d); d = NULL;
output_ptr = (float*)malloc(length);
img_format.image_channel_order = CL_RGBA;
img_format.image_channel_data_type = CL_FLOAT;
streams[0] = create_image_3d(context, CL_MEM_READ_ONLY, &img_format, img_width, img_height, img_depth, 0, 0, NULL, &err);
test_error(err, "create_image_3d failed");
streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
test_error(err, "clCreateBuffer failed");
err = clEnqueueWriteImage(queue, streams[0], CL_TRUE, origin, region, 0, 0, input_ptr, 0, NULL, NULL);
test_error(err, "clEnqueueWriteImage failed");
err = create_single_kernel_helper(context, &program, &kernel, 1, &rgbaFFFF_kernel_code, "test_rgbaFFFF" );
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 streams[0], &streams[0]);
err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]);
err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler);
test_error(err, "clSetKernelArg failed");
threads[0] = (unsigned int)img_width;
threads[1] = (unsigned int)img_height;
threads[2] = (unsigned int)img_depth;
err = clEnqueueNDRangeKernel(queue, kernel, 3, NULL, threads, NULL, 0, NULL, NULL);
test_error(err, "clEnqueueNDRangeKernel failed");
err = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
test_error(err, "clEnqueueReadBuffer failed");
err = verify_float_image(input_ptr, output_ptr, img_width, img_height, img_depth);
// cleanup
clReleaseSampler(sampler);
clReleaseMemObject(streams[0]);
clReleaseMemObject(streams[1]);
clReleaseKernel(kernel);
clReleaseProgram(program);
free(input_ptr);
free(output_ptr);
return err;
}

146
test_conformance/compatibility/test_conformance/basic/test_readimage3d_int16.c
View File

@@ -1,146 +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 <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#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<w*h*d*4; i++)
ptr[i] = (cl_ushort)genrand_int32(data);
return ptr;
}
static int
verify_16bit_image(cl_ushort *image, cl_ushort *outptr, int w, int h, int d)
{
int i;
for (i=0; i<w*h*d*4; i++)
{
if (outptr[i] != image[i])
{
log_error("READ_IMAGE3D_RGBA_UNORM_INT16 test failed\n");
return -1;
}
}
log_info("READ_IMAGE3D_RGBA_UNORM_INT16 test passed\n");
return 0;
}
int test_readimage3d_int16(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
cl_mem streams[2];
cl_program program;
cl_kernel kernel;
cl_image_format img_format;
cl_ushort *input_ptr, *output_ptr;
size_t threads[3];
int img_width = 64;
int img_height = 64;
int img_depth = 64;
int err;
size_t origin[3] = {0, 0, 0};
size_t region[3] = {img_width, img_height, img_depth};
size_t length = img_width * img_height * img_depth * 4 * sizeof(cl_ushort);
PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( device )
MTdata d = init_genrand( gRandomSeed );
input_ptr = generate_16bit_image(img_width, img_height, img_depth, d);
free_mtdata(d); d = NULL;
output_ptr = (cl_ushort*)malloc(length);
img_format.image_channel_order = CL_RGBA;
img_format.image_channel_data_type = CL_UNORM_INT16;
streams[0] = create_image_3d(context, CL_MEM_READ_ONLY, &img_format, img_width, img_height, img_depth, 0, 0, NULL, &err);
test_error(err, "create_image_3d failed");
streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
test_error(err, "clCreateBuffer failed");
err = clEnqueueWriteImage(queue, streams[0], CL_TRUE, origin, region, 0, 0, input_ptr, 0, NULL, NULL);
test_error(err, "clEnqueueWriteImage failed");
err = create_single_kernel_helper(context, &program, &kernel, 1, &rgba16_kernel_code, "test_rgba16" );
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 streams[0], &streams[0]);
err |= clSetKernelArg(kernel, 1, sizeof streams[1], &streams[1]);
err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler);
test_error(err, "clSetKernelArg failed");
threads[0] = (unsigned int)img_width;
threads[1] = (unsigned int)img_height;
threads[2] = (unsigned int)img_depth;
err = clEnqueueNDRangeKernel(queue, kernel, 3, NULL, threads, NULL, 0, NULL, NULL);
test_error(err, "clEnqueueNDRangeKernel failed");
err = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
test_error(err, "clEnqueueReadBuffer failed");
err = verify_16bit_image(input_ptr, output_ptr, img_width, img_height, img_depth);
// cleanup
clReleaseSampler(sampler);
clReleaseMemObject(streams[0]);
clReleaseMemObject(streams[1]);
clReleaseKernel(kernel);
clReleaseProgram(program);
free(input_ptr);
free(output_ptr);
return err;
}

167
test_conformance/compatibility/test_conformance/basic/test_readimage_fp32.c
View File

@@ -1,167 +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 <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#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<w*h*4; i++)
ptr[i] = get_random_float(-0x40000000, 0x40000000, d);
return ptr;
}
static int
verify_float_image(float *image, float *outptr, int w, int h)
{
int i;
for (i=0; i<w*h*4; i++)
{
if (outptr[i] != image[i])
{
log_error("READ_IMAGE_RGBA_FLOAT test failed\n");
return -1;
}
}
log_info("READ_IMAGE_RGBA_FLOAT test passed\n");
return 0;
}
int test_readimage_fp32(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
cl_mem streams[2];
cl_program program;
cl_kernel kernel;
cl_image_format img_format;
float *input_ptr, *output_ptr;
size_t threads[2];
int img_width = 512;
int img_height = 512;
int 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 )
d = init_genrand( gRandomSeed );
input_ptr = generate_float_image(img_width, img_height, d);
free_mtdata(d); d = NULL;
output_ptr = (float*)malloc(length);
img_format.image_channel_order = CL_RGBA;
img_format.image_channel_data_type = CL_FLOAT;
streams[0] = create_image_2d(context, 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;
}
streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
if (!streams[1])
{
log_error("clCreateArray failed\n");
return -1;
}
err = clEnqueueWriteImage(queue, streams[0], CL_TRUE, origin, region, 0, 0, input_ptr, 0, NULL, NULL);
if (err != CL_SUCCESS)
{
log_error("clWriteImage failed\n");
return -1;
}
err = create_single_kernel_helper(context, &program, &kernel, 1, &rgbaFFFF_kernel_code, "test_rgbaFFFF" );
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 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");
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("%s clEnqueueNDRangeKernel failed\n", __FUNCTION__);
return -1;
}
err = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
if (err != CL_SUCCESS)
{
log_error("clEnqueueReadBuffer failed\n");
return -1;
}
err = verify_float_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);
return err;
}

166
test_conformance/compatibility/test_conformance/basic/test_readimage_int16.c
View File

@@ -1,166 +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 <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#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<w*h*4; i++)
ptr[i] = (cl_ushort)genrand_int32(d);
return ptr;
}
static int
verify_16bit_image(cl_ushort *image, cl_ushort *outptr, int w, int h)
{
int i;
for (i=0; i<w*h*4; i++)
{
if (outptr[i] != image[i])
{
log_error("READ_IMAGE_RGBA_UNORM_INT16 test failed\n");
return -1;
}
}
log_info("READ_IMAGE_RGBA_UNORM_INT16 test passed\n");
return 0;
}
int test_readimage_int16(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
cl_mem streams[2];
cl_program program;
cl_kernel kernel;
cl_image_format img_format;
cl_ushort *input_ptr, *output_ptr;
size_t threads[2];
int img_width = 512;
int img_height = 512;
int 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(cl_ushort);
MTdata d;
PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
d = init_genrand( gRandomSeed );
input_ptr = generate_16bit_image(img_width, img_height, d);
free_mtdata(d); d = NULL;
output_ptr = (cl_ushort*)malloc(length);
img_format.image_channel_order = CL_RGBA;
img_format.image_channel_data_type = CL_UNORM_INT16;
streams[0] = create_image_2d(context, 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;
}
streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
if (!streams[1])
{
log_error("clCreateArray failed\n");
return -1;
}
err = clEnqueueWriteImage(queue, streams[0], CL_TRUE, origin, region, 0, 0, input_ptr, 0, NULL, NULL);
if (err != CL_SUCCESS)
{
log_error("clWriteImage failed\n");
return -1;
}
err = create_single_kernel_helper(context, &program, &kernel, 1, &rgba16_kernel_code, "test_rgba16" );
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 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");
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("%s clEnqueueNDRangeKernel failed\n", __FUNCTION__);
return -1;
}
err = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
if (err != CL_SUCCESS)
{
log_error("clEnqueueReadBuffer failed\n");
return -1;
}
err = verify_16bit_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);
return err;
}