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

Allow CL_HALF_FLOAT denorm flushing for write tests (#452) (#453)

Browse Source 
* Allow CL_HALF_FLOAT denorm flushing for write tests (#452)

* On mismatch, add relaxation when denormal half result is expected
* Refactor to use common validation function
* Clean up some diagnostics

* Fix review comments

- use cl_half
- remove extraneous casts
- replace literals with sizeof()

* Document rollover trick for IsHalfSubnormal
This commit is contained in:
Jim Lewis
2020-04-16 02:23:44 -07:00
committed by GitHub
parent 2fa8611862
commit f3a3ec2b47
9 changed files with 95 additions and 134 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
test_common/harness/conversions.h
View File

@@ -117,6 +117,11 @@ static inline int IsDoubleSubnormal( double x )
#endif
}
static inline int IsHalfSubnormal( cl_half x )
{
return ( ( x & 0x7fffU ) - 1U ) < 0x03ffU;
}
#if defined(__cplusplus)
}
#endif

9
test_common/harness/imageHelpers.h
View File

@@ -640,7 +640,14 @@ protected:
extern int DetectFloatToHalfRoundingMode( cl_command_queue ); // Returns CL_SUCCESS on success
int inline is_half_nan( cl_ushort half ){ return (half & 0x7fff) > 0x7c00; }
// sign bit: don't care, exponent: maximum value, significand: non-zero
static int inline is_half_nan( cl_ushort half ){ return ( half & 0x7fff ) > 0x7c00; }
// sign bit: don't care, exponent: zero, significand: non-zero
static int inline is_half_denorm( cl_ushort half ){ return IsHalfSubnormal( half ); }
// sign bit: don't care, exponent: zero, significand: zero
static int inline is_half_zero( cl_ushort half ){ return ( half & 0x7fff ) == 0; }
cl_ushort convert_float_to_half( cl_float f );
cl_float convert_half_to_float( cl_ushort h );

3
test_conformance/half/cl_utils.h
View File

@@ -114,7 +114,8 @@ static inline cl_ulong DoubleFromUInt( cl_uint bits )
static inline int IsHalfSubnormal( uint16_t x )
{
return ((x&0x7fffU)-1U) < 0x03ffU;
// this relies on interger overflow to exclude 0 as a subnormal
return ( ( x & 0x7fffU ) - 1U ) < 0x03ffU;
}
// prevent silent failures due to missing FLT_RADIX

21
test_conformance/images/kernel_read_write/test_iterations.cpp
View File

@@ -1155,7 +1155,7 @@ bool validate_float_write_results( float *expected, float *actual, image_descrip
{
bool pass = true;
// Compare floats
if( memcmp( expected, actual, 4 * get_format_channel_count( imageInfo->format ) ) != 0 )
if( memcmp( expected, actual, sizeof( cl_float ) * get_format_channel_count( imageInfo->format ) ) != 0 )
{
// 8.3.3 Fix up cases where we have NaNs or flushed denorms; "all other values must be preserved"
for ( size_t j = 0; j < get_format_channel_count( imageInfo->format ); j++ )
@@ -1171,6 +1171,25 @@ bool validate_float_write_results( float *expected, float *actual, image_descrip
return pass;
}
bool validate_half_write_results( cl_half *expected, cl_half *actual, image_descriptor *imageInfo )
{
bool pass = true;
// Compare half floats
if (memcmp(expected, actual, sizeof( cl_half ) * get_format_channel_count(imageInfo->format)) != 0) {
// 8.3.2 Fix up cases where we have NaNs or generated half denormals
for ( size_t j = 0; j < get_format_channel_count( imageInfo->format ); j++ ) {
if ( is_half_nan( expected[j] ) && is_half_nan( actual[j] ) )
continue;
if ( is_half_denorm( expected[j] ) && is_half_zero( actual[j] ) )
continue;
pass = false;
break;
}
}
return pass;
}
int test_read_image_2D( cl_context context, cl_command_queue queue, cl_kernel kernel,
image_descriptor *imageInfo, image_sampler_data *imageSampler,
bool useFloatCoords, ExplicitType outputType, MTdata d )

42
test_conformance/images/kernel_read_write/test_write_1D.cpp
View File

@@ -27,6 +27,7 @@ extern int gtestTypesToRun;
extern bool gDeviceLt20;
extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo );
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"
@@ -376,37 +377,22 @@ int test_write_image_1D( cl_device_id device, cl_context context, cl_command_que
}
else if( imageInfo->format->image_channel_data_type == CL_HALF_FLOAT )
{
// Compare half floats
if( memcmp( resultBuffer, resultPtr, 2 * get_format_channel_count( imageInfo->format ) ) != 0 )
cl_half *e = (cl_half *)resultBuffer;
cl_half *a = (cl_half *)resultPtr;
if( !validate_half_write_results( e, a, imageInfo ) )
{
cl_ushort *e = (cl_ushort *)resultBuffer;
cl_ushort *a = (cl_ushort *)resultPtr;
int err_cnt = 0;
//Fix up cases where we have NaNs
for( size_t j = 0; j < get_format_channel_count( imageInfo->format ); j++ )
totalErrors++;
log_error( "ERROR: Sample %ld did not validate! (%s)\n", i, mem_flag_names[ mem_flag_index ] );
log_error( " Expected: 0x%04x 0x%04x 0x%04x 0x%04x\n", e[ 0 ], e[ 1 ], e[ 2 ], e[ 3 ] );
log_error( " Actual: 0x%04x 0x%04x 0x%04x 0x%04x\n", a[ 0 ], a[ 1 ], a[ 2 ], a[ 3 ] );
if( inputType == kFloat )
{
if( is_half_nan( e[j] ) && is_half_nan(a[j]) )
continue;
if( e[j] != a[j] )
err_cnt++;
}
if( err_cnt )
{
totalErrors++;
log_error( "ERROR: Sample %ld (%ld) did not validate! (%s)\n", i, x, mem_flag_names[mem_flag_index] );
log_error( " Expected: 0x%04x 0x%04x 0x%04x 0x%04x\n", e[0], e[1], e[2], e[3] );
log_error( " Actual: 0x%04x 0x%04x 0x%04x 0x%04x\n", a[0], a[1], a[2], a[3] );
if( inputType == kFloat )
{
float *p = (float *)(char *)imagePtr;
log_error( " Source: %a %a %a %a\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
log_error( " : %12.24f %12.24f %12.24f %12.24f\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
}
if( ( --numTries ) == 0 )
return 1;
float *p = (float *)imagePtr;
log_error( " Source: %a %a %a %a\n", p[ 0 ], p[ 1 ], p[ 2] , p[ 3] );
log_error( " : %12.24f %12.24f %12.24f %12.24f\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
}
if( ( --numTries ) == 0 )
return 1;
}
}
else

45
test_conformance/images/kernel_read_write/test_write_1D_array.cpp
View File

@@ -26,6 +26,7 @@ extern int gtestTypesToRun;
extern bool gDeviceLt20;
extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo );
extern bool validate_half_write_results( cl_half *expected, cl_half *actual, image_descriptor *imageInfo );
const char *readwrite1DArrayKernelSourcePattern =
"__kernel void sample_kernel( __global %s4 *input, read_write image1d_array_t output %s)\n"
@@ -395,38 +396,22 @@ int test_write_image_1D_array( cl_device_id device, cl_context context, cl_comma
}
else if( imageInfo->format->image_channel_data_type == CL_HALF_FLOAT )
{
// Compare half floats
if( memcmp( resultBuffer, resultPtr, 2 * get_format_channel_count( imageInfo->format ) ) != 0 )
cl_half *e = (cl_half *)resultBuffer;
cl_half *a = (cl_half *)resultPtr;
if( !validate_half_write_results( e, a, imageInfo ) )
{
cl_ushort *e = (cl_ushort *)resultBuffer;
cl_ushort *a = (cl_ushort *)resultPtr;
int err_cnt = 0;
//Fix up cases where we have NaNs
for( size_t j = 0; j < get_format_channel_count( imageInfo->format ); j++ )
totalErrors++;
log_error( "ERROR: Sample %ld (%ld,%ld) did not validate! (%s)\n", i, x, y, mem_flag_names[ mem_flag_index ] );
log_error( " Expected: 0x%04x 0x%04x 0x%04x 0x%04x\n", e[ 0 ], e[ 1 ], e[ 2 ], e[ 3 ] );
log_error( " Actual: 0x%04x 0x%04x 0x%04x 0x%04x\n", a[ 0 ], a[ 1 ], a[ 2 ], a[ 3 ] );
if( inputType == kFloat )
{
if( is_half_nan( e[j] ) && is_half_nan(a[j]) )
continue;
if( e[j] != a[j] )
err_cnt++;
}
if( err_cnt )
{
totalErrors++;
log_error( "ERROR: Sample %ld (%ld,%ld) did not validate! (%s)\n", i, x, y, mem_flag_names[mem_flag_index] );
log_error( " Expected: 0x%04x 0x%04x 0x%04x 0x%04x\n", e[0], e[1], e[2], e[3] );
log_error( " Actual: 0x%04x 0x%04x 0x%04x 0x%04x\n", a[0], a[1], a[2], a[3] );
if( inputType == kFloat )
{
float *p = (float *)(char *)imagePtr;
log_error( " Source: %a %a %a %a\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
log_error( " : %12.24f %12.24f %12.24f %12.24f\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
}
if( ( --numTries ) == 0 )
return 1;
float *p = (float *)imagePtr;
log_error( " Source: %a %a %a %a\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
log_error( " : %12.24f %12.24f %12.24f %12.24f\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
}
if( ( --numTries ) == 0 )
return 1;
}
}
else
@@ -522,7 +507,7 @@ int test_write_image_1D_array( cl_device_id device, cl_context context, cl_comma
break;
}
float *v = (float *)(char *)imagePtr;
float *v = (float *)imagePtr;
log_error( " src: %g %g %g %g\n", v[ 0 ], v[ 1], v[ 2 ], v[ 3 ] );
log_error( " : %a %a %a %a\n", v[ 0 ], v[ 1], v[ 2 ], v[ 3 ] );
log_error( " src: %12.24f %12.24f %12.24f %12.24f\n", v[0 ], v[ 1], v[ 2 ], v[ 3 ] );

30
test_conformance/images/kernel_read_write/test_write_2D_array.cpp
View File

@@ -27,6 +27,7 @@ extern int gtestTypesToRun;
extern bool gDeviceLt20;
extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo );
extern bool validate_half_write_results( cl_half *expected, cl_half *actual, image_descriptor *imageInfo );
// Utility function to clamp down image sizes for certain tests to avoid
// using too much memory.
@@ -416,40 +417,25 @@ int test_write_image_2D_array( cl_device_id device, cl_context context, cl_comma
}
else if( imageInfo->format->image_channel_data_type == CL_HALF_FLOAT )
{
// Compare half floats
if( memcmp( resultBuffer, resultPtr, 2 * get_format_channel_count( imageInfo->format ) ) != 0 )
cl_half *e = (cl_half *)resultBuffer;
cl_half *a = (cl_half *)resultPtr;
if( !validate_half_write_results( e, a, imageInfo ) )
{
cl_ushort *e = (cl_ushort *)resultBuffer;
cl_ushort *a = (cl_ushort *)resultPtr;
int err_cnt = 0;
//Fix up cases where we have NaNs
for( size_t j = 0; j < get_format_channel_count( imageInfo->format ); j++ )
{
if( is_half_nan( e[j] ) && is_half_nan(a[j]) )
continue;
if( e[j] != a[j] )
err_cnt++;
}
if( err_cnt )
{
totalErrors++;
log_error( "ERROR: Sample %ld (%ld,%ld) did not validate! (%s)\n", i, x, y, mem_flag_names[mem_flag_index] );
log_error( "ERROR: Sample %ld (%ld,%ld,%ld) did not validate! (%s)\n", i, x, y, z, mem_flag_names[ mem_flag_index ] );
unsigned short *e = (unsigned short *)resultBuffer;
unsigned short *a = (unsigned short *)resultPtr;
log_error( " Expected: 0x%04x 0x%04x 0x%04x 0x%04x\n", e[0], e[1], e[2], e[3] );
log_error( " Actual: 0x%04x 0x%04x 0x%04x 0x%04x\n", a[0], a[1], a[2], a[3] );
log_error( " Expected: 0x%04x 0x%04x 0x%04x 0x%04x\n", e[ 0 ], e[ 1 ], e[ 2 ], e[ 3 ] );
log_error( " Actual: 0x%04x 0x%04x 0x%04x 0x%04x\n", a[ 0 ], a[ 1 ], a[ 2 ], a[ 3 ] );
if( inputType == kFloat )
{
float *p = (float *)(char *)imagePtr;
float *p = (float *)imagePtr;
log_error( " Source: %a %a %a %a\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
log_error( " : %12.24f %12.24f %12.24f %12.24f\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
}
if( ( --numTries ) == 0 )
return 1;
}
}
}
else
{

30
test_conformance/images/kernel_read_write/test_write_3D.cpp
View File

@@ -27,6 +27,7 @@ extern int gtestTypesToRun;
extern bool gDeviceLt20;
extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo );
extern bool validate_half_write_results( cl_half *expected, cl_half *actual, image_descriptor *imageInfo );
// Utility function to clamp down image sizes for certain tests to avoid
// using too much memory.
@@ -423,40 +424,25 @@ int test_write_image_3D( cl_device_id device, cl_context context, cl_command_que
}
else if( imageInfo->format->image_channel_data_type == CL_HALF_FLOAT )
{
// Compare half floats
if( memcmp( resultBuffer, resultPtr, 2 * get_format_channel_count( imageInfo->format ) ) != 0 )
cl_half *e = (cl_half *)resultBuffer;
cl_half *a = (cl_half *)resultPtr;
if( !validate_half_write_results( e, a, imageInfo ) )
{
cl_ushort *e = (cl_ushort *)resultBuffer;
cl_ushort *a = (cl_ushort *)resultPtr;
int err_cnt = 0;
//Fix up cases where we have NaNs
for( size_t j = 0; j < get_format_channel_count( imageInfo->format ); j++ )
{
if( is_half_nan( e[j] ) && is_half_nan(a[j]) )
continue;
if( e[j] != a[j] )
err_cnt++;
}
if( err_cnt )
{
totalErrors++;
log_error( "ERROR: Sample %ld (%ld,%ld) did not validate! (%s)\n", i, x, y, mem_flag_names[mem_flag_index] );
log_error( "ERROR: Sample %ld (%ld,%ld,%ld) did not validate! (%s)\n", i, x, y, z, mem_flag_names[ mem_flag_index ] );
unsigned short *e = (unsigned short *)resultBuffer;
unsigned short *a = (unsigned short *)resultPtr;
log_error( " Expected: 0x%04x 0x%04x 0x%04x 0x%04x\n", e[0], e[1], e[2], e[3] );
log_error( " Actual: 0x%04x 0x%04x 0x%04x 0x%04x\n", a[0], a[1], a[2], a[3] );
log_error( " Expected: 0x%04x 0x%04x 0x%04x 0x%04x\n", e[ 0 ], e[ 1 ], e[ 2 ], e[ 3 ] );
log_error( " Actual: 0x%04x 0x%04x 0x%04x 0x%04x\n", a[ 0 ], a[ 1 ], a[ 2 ], a[ 3 ] );
if( inputType == kFloat )
{
float *p = (float *)(char *)imagePtr;
float *p = (float *)imagePtr;
log_error( " Source: %a %a %a %a\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
log_error( " : %12.24f %12.24f %12.24f %12.24f\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
}
if( ( --numTries ) == 0 )
return 1;
}
}
}
else
{

44
test_conformance/images/kernel_read_write/test_write_image.cpp
View File

@@ -29,7 +29,9 @@ extern int test_write_image_1D_set( cl_device_id device, cl_context context, cl_
extern int test_write_image_3D_set( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, ExplicitType inputType, MTdata d );
extern int test_write_image_1D_array_set( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, ExplicitType inputType, MTdata d );
extern int test_write_image_2D_array_set( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, ExplicitType inputType, MTdata d );
extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo );
extern bool validate_half_write_results( cl_half *expected, cl_half *actual, image_descriptor *imageInfo );
const char *writeKernelSourcePattern =
"__kernel void sample_kernel( __global %s%s *input, write_only %s output %s)\n"
@@ -440,38 +442,22 @@ int test_write_image( cl_device_id device, cl_context context, cl_command_queue
}
else if( imageInfo->format->image_channel_data_type == CL_HALF_FLOAT )
{
// Compare half floats
if( memcmp( resultBuffer, resultPtr, 2 * get_format_channel_count( imageInfo->format ) ) != 0 )
cl_half *e = (cl_half *)resultBuffer;
cl_half *a = (cl_half *)resultPtr;
if( !validate_half_write_results( e, a, imageInfo ) )
{
cl_ushort *e = (cl_ushort *)resultBuffer;
cl_ushort *a = (cl_ushort *)resultPtr;
int err_cnt = 0;
//Fix up cases where we have NaNs
for( size_t j = 0; j < get_format_channel_count( imageInfo->format ); j++ )
totalErrors++;
log_error( "ERROR: Sample %ld (%ld,%ld) did not validate! (%s)\n", i, x, y, mem_flag_names[ mem_flag_index ] );
log_error( " Expected: 0x%04x 0x%04x 0x%04x 0x%04x\n", e[ 0 ], e[ 1 ], e[ 2 ], e[ 3 ] );
log_error( " Actual: 0x%04x 0x%04x 0x%04x 0x%04x\n", a[ 0 ], a[ 1 ], a[ 2 ], a[ 3 ] );
if( inputType == kFloat )
{
if( is_half_nan( e[j] ) && is_half_nan(a[j]) )
continue;
if( e[j] != a[j] )
err_cnt++;
}
if( err_cnt )
{
totalErrors++;
log_error( "ERROR: Sample %ld (%ld,%ld) did not validate! (%s)\n", i, x, y, mem_flag_names[mem_flag_index] );
log_error( " Expected: 0x%04x 0x%04x 0x%04x 0x%04x\n", e[0], e[1], e[2], e[3] );
log_error( " Actual: 0x%04x 0x%04x 0x%04x 0x%04x\n", a[0], a[1], a[2], a[3] );
if( inputType == kFloat )
{
float *p = (float *)(char *)imagePtr;
log_error( " Source: %a %a %a %a\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
log_error( " : %12.24f %12.24f %12.24f %12.24f\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
}
if( ( --numTries ) == 0 )
return 1;
float *p = (float *)(char *)imagePtr;
log_error( " Source: %a %a %a %a\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
log_error( " : %12.24f %12.24f %12.24f %12.24f\n", p[ 0 ], p[ 1 ], p[ 2 ], p[ 3 ] );
}
if( ( --numTries ) == 0 )
return 1;
}
}
else