diff --git a/test_common/harness/conversions.h b/test_common/harness/conversions.h index aa3cb6b4..a9dd6dcf 100644 --- a/test_common/harness/conversions.h +++ b/test_common/harness/conversions.h @@ -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 diff --git a/test_common/harness/imageHelpers.h b/test_common/harness/imageHelpers.h index 1ac1faca..26c97603 100644 --- a/test_common/harness/imageHelpers.h +++ b/test_common/harness/imageHelpers.h @@ -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 ); diff --git a/test_conformance/half/cl_utils.h b/test_conformance/half/cl_utils.h index 62869324..814cb902 100644 --- a/test_conformance/half/cl_utils.h +++ b/test_conformance/half/cl_utils.h @@ -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 diff --git a/test_conformance/images/kernel_read_write/test_iterations.cpp b/test_conformance/images/kernel_read_write/test_iterations.cpp index b0554509..050a805d 100644 --- a/test_conformance/images/kernel_read_write/test_iterations.cpp +++ b/test_conformance/images/kernel_read_write/test_iterations.cpp @@ -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 ) 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 86592c4c..ca022629 100644 --- a/test_conformance/images/kernel_read_write/test_write_1D.cpp +++ b/test_conformance/images/kernel_read_write/test_write_1D.cpp @@ -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 diff --git a/test_conformance/images/kernel_read_write/test_write_1D_array.cpp b/test_conformance/images/kernel_read_write/test_write_1D_array.cpp index 62b688c4..b91bf1cf 100644 --- a/test_conformance/images/kernel_read_write/test_write_1D_array.cpp +++ b/test_conformance/images/kernel_read_write/test_write_1D_array.cpp @@ -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 ] ); diff --git a/test_conformance/images/kernel_read_write/test_write_2D_array.cpp b/test_conformance/images/kernel_read_write/test_write_2D_array.cpp index cd1c8495..4524c6cd 100644 --- a/test_conformance/images/kernel_read_write/test_write_2D_array.cpp +++ b/test_conformance/images/kernel_read_write/test_write_2D_array.cpp @@ -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 { diff --git a/test_conformance/images/kernel_read_write/test_write_3D.cpp b/test_conformance/images/kernel_read_write/test_write_3D.cpp index 2dc818ca..7440bd6e 100644 --- a/test_conformance/images/kernel_read_write/test_write_3D.cpp +++ b/test_conformance/images/kernel_read_write/test_write_3D.cpp @@ -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 { diff --git a/test_conformance/images/kernel_read_write/test_write_image.cpp b/test_conformance/images/kernel_read_write/test_write_image.cpp index 8d231445..f6d9235c 100644 --- a/test_conformance/images/kernel_read_write/test_write_image.cpp +++ b/test_conformance/images/kernel_read_write/test_write_image.cpp @@ -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