aboutsummaryrefslogtreecommitdiff
path: root/test_conformance/images/kernel_read_write/test_write_1D.cpp
diff options
context:
space:
mode:
Diffstat (limited to 'test_conformance/images/kernel_read_write/test_write_1D.cpp')
-rw-r--r--test_conformance/images/kernel_read_write/test_write_1D.cpp112
1 files changed, 79 insertions, 33 deletions
diff --git a/test_conformance/images/kernel_read_write/test_write_1D.cpp b/test_conformance/images/kernel_read_write/test_write_1D.cpp
index 41983edf..5f726796 100644
--- a/test_conformance/images/kernel_read_write/test_write_1D.cpp
+++ b/test_conformance/images/kernel_read_write/test_write_1D.cpp
@@ -14,6 +14,7 @@
// limitations under the License.
//
#include "../testBase.h"
+#include "test_common.h"
#if !defined(_WIN32)
#include <sys/mman.h>
@@ -26,20 +27,24 @@ extern bool validate_float_write_results( float *expected, float *actual, image_
extern bool validate_half_write_results( cl_half *expected, cl_half *actual, image_descriptor* imageInfo );
const char *readwrite1DKernelSourcePattern =
-"__kernel void sample_kernel( __global %s4 *input, read_write image1d_t output %s)\n"
-"{\n"
-" int tidX = get_global_id(0);\n"
-" int offset = tidX;\n"
-" write_image%s( output, tidX %s, input[ offset ]);\n"
-"}";
+ "%s\n"
+ "__kernel void sample_kernel( __global %s4 *input, read_write image1d_t "
+ "output %s)\n"
+ "{\n"
+ " int tidX = get_global_id(0);\n"
+ " int offset = tidX;\n"
+ " write_image%s( output, tidX %s, input[ offset ]);\n"
+ "}";
const char *write1DKernelSourcePattern =
-"__kernel void sample_kernel( __global %s4 *input, write_only image1d_t output %s)\n"
-"{\n"
-" int tidX = get_global_id(0);\n"
-" int offset = tidX;\n"
-" write_image%s( output, tidX %s, input[ offset ]);\n"
-"}";
+ "%s\n"
+ "__kernel void sample_kernel( __global %s4 *input, write_only image1d_t "
+ "output %s)\n"
+ "{\n"
+ " int tidX = get_global_id(0);\n"
+ " int offset = tidX;\n"
+ " write_image%s( output, tidX %s, input[ offset ]);\n"
+ "}";
int test_write_image_1D( cl_device_id device, cl_context context, cl_command_queue queue, cl_kernel kernel,
image_descriptor *imageInfo, ExplicitType inputType, MTdata d )
@@ -395,6 +400,8 @@ int test_write_image_1D( cl_device_id device, cl_context context, cl_command_que
}
else
{
+ filter_undefined_bits(imageInfo, resultPtr);
+
// Exact result passes every time
if( memcmp( resultBuffer, resultPtr, get_pixel_size( imageInfo->format ) ) != 0 )
{
@@ -403,21 +410,8 @@ int test_write_image_1D( cl_device_id device, cl_context context, cl_command_que
float errors[4] = {NAN, NAN, NAN, NAN};
pack_image_pixel_error( (float *)imagePtr, imageInfo->format, resultBuffer, errors );
- // We are allowed 0.6 absolute error vs. infinitely precise for some normalized formats
- if( 0 == forceCorrectlyRoundedWrites &&
- (
- imageInfo->format->image_channel_data_type == CL_UNORM_INT8 ||
- imageInfo->format->image_channel_data_type == CL_UNORM_INT_101010 ||
- imageInfo->format->image_channel_data_type == CL_UNORM_INT16 ||
- imageInfo->format->image_channel_data_type == CL_SNORM_INT8 ||
- imageInfo->format->image_channel_data_type == CL_SNORM_INT16
- ))
- {
- if( ! (fabsf( errors[0] ) > 0.6f) && ! (fabsf( errors[1] ) > 0.6f) &&
- ! (fabsf( errors[2] ) > 0.6f) && ! (fabsf( errors[3] ) > 0.6f) )
- failure = 0;
- }
-
+ failure = filter_rounding_errors(
+ forceCorrectlyRoundedWrites, imageInfo, errors);
if( failure )
{
@@ -458,6 +452,56 @@ int test_write_image_1D( cl_device_id device, cl_context context, cl_command_que
log_error( " Actual: 0x%2.2x 0x%2.2x 0x%2.2x 0x%2.2x\n", ((cl_uchar*)resultPtr)[0], ((cl_uchar*)resultPtr)[1], ((cl_uchar*)resultPtr)[2], ((cl_uchar*)resultPtr)[3] );
log_error( " Error: %f %f %f %f\n", errors[0], errors[1], errors[2], errors[3] );
break;
+ case CL_UNORM_SHORT_565: {
+ cl_uint *ref_value =
+ (cl_uint *)resultBuffer;
+ cl_uint *test_value =
+ (cl_uint *)resultPtr;
+
+ log_error(" Expected: 0x%2.2x Actual: "
+ "0x%2.2x \n",
+ ref_value[0], test_value[0]);
+
+ log_error(" Expected: 0x%2.2x "
+ "0x%2.2x 0x%2.2x \n",
+ ref_value[0] & 0x1F,
+ (ref_value[0] >> 5) & 0x3F,
+ (ref_value[0] >> 11) & 0x1F);
+ log_error(" Actual: 0x%2.2x "
+ "0x%2.2x 0x%2.2x \n",
+ test_value[0] & 0x1F,
+ (test_value[0] >> 5) & 0x3F,
+ (test_value[0] >> 11) & 0x1F);
+ log_error(" Error: %f %f %f %f\n",
+ errors[0], errors[1],
+ errors[2]);
+ break;
+ }
+ case CL_UNORM_SHORT_555: {
+ cl_uint *ref_value =
+ (cl_uint *)resultBuffer;
+ cl_uint *test_value =
+ (cl_uint *)resultPtr;
+
+ log_error(" Expected: 0x%2.2x Actual: "
+ "0x%2.2x \n",
+ ref_value[0], test_value[0]);
+
+ log_error(" Expected: 0x%2.2x "
+ "0x%2.2x 0x%2.2x \n",
+ ref_value[0] & 0x1F,
+ (ref_value[0] >> 5) & 0x1F,
+ (ref_value[0] >> 10) & 0x1F);
+ log_error(" Actual: 0x%2.2x "
+ "0x%2.2x 0x%2.2x \n",
+ test_value[0] & 0x1F,
+ (test_value[0] >> 5) & 0x1F,
+ (test_value[0] >> 10) & 0x1F);
+ log_error(" Error: %f %f %f %f\n",
+ errors[0], errors[1],
+ errors[2]);
+ break;
+ }
case CL_UNORM_INT16:
case CL_SNORM_INT16:
case CL_UNSIGNED_INT16:
@@ -574,12 +618,14 @@ int test_write_image_1D_set(cl_device_id device, cl_context context,
KernelSourcePattern = readwrite1DKernelSourcePattern;
}
- sprintf( programSrc,
- KernelSourcePattern,
- get_explicit_type_name( inputType ),
- gTestMipmaps ? ", int lod" : "",
- readFormat,
- gTestMipmaps ? ", lod" :"" );
+ sprintf(
+ programSrc, KernelSourcePattern,
+ gTestMipmaps
+ ? "#pragma OPENCL EXTENSION cl_khr_mipmap_image: enable\n#pragma "
+ "OPENCL EXTENSION cl_khr_mipmap_image_writes: enable"
+ : "",
+ get_explicit_type_name(inputType), gTestMipmaps ? ", int lod" : "",
+ readFormat, gTestMipmaps ? ", lod" : "");
ptr = programSrc;
error = create_single_kernel_helper(context, &program, &kernel, 1, &ptr,
generated by cgit v1.2.3 (git 2.25.1) at 2024-06-10 06:37:19 +0000