From 040321d8b90a95866e329e3be2a42aa151d1b5a1 Mon Sep 17 00:00:00 2001 From: Jim Lewis Date: Fri, 15 Nov 2019 07:38:15 -0600 Subject: [PATCH] Allow CL_FLOAT denorm flushing for write tests (#28) (#456) * Require exact for match normals, instead of arbitrary .005 relative error * Add relaxation to allow 0 when float denormal is expected * Refactor to use common validation function --- test_common/harness/imageHelpers.cpp | 4 ++-- .../kernel_read_write/test_iterations.cpp | 20 ++++++++++++++++++ .../kernel_read_write/test_write_1D.cpp | 21 +++++++------------ .../kernel_read_write/test_write_1D_array.cpp | 18 ++++++---------- .../kernel_read_write/test_write_2D_array.cpp | 19 +++++------------ .../kernel_read_write/test_write_3D.cpp | 19 +++++------------ .../kernel_read_write/test_write_image.cpp | 18 +++++----------- 7 files changed, 50 insertions(+), 69 deletions(-) diff --git a/test_common/harness/imageHelpers.cpp b/test_common/harness/imageHelpers.cpp index 85e28499..7ad99226 100644 --- a/test_common/harness/imageHelpers.cpp +++ b/test_common/harness/imageHelpers.cpp @@ -1808,7 +1808,7 @@ static inline void check_for_denorms(float a[4], int *containsDenorms ) { for( int i = 0; i < 4; i++ ) { - if( fabsf(a[i]) < FLT_MIN ) + if( IsFloatSubnormal( a[i] ) ) a[i] = copysignf( 0.0f, a[i] ); } } @@ -1816,7 +1816,7 @@ static inline void check_for_denorms(float a[4], int *containsDenorms ) { for( int i = 0; i < 4; i++ ) { - if( fabs(a[i]) < FLT_MIN ) + if( IsFloatSubnormal( a[i] ) ) { *containsDenorms = 1; break; diff --git a/test_conformance/images/kernel_read_write/test_iterations.cpp b/test_conformance/images/kernel_read_write/test_iterations.cpp index afb879fb..6d064af8 100644 --- a/test_conformance/images/kernel_read_write/test_iterations.cpp +++ b/test_conformance/images/kernel_read_write/test_iterations.cpp @@ -1151,6 +1151,26 @@ int validate_image_2D_sRGB_results(void *imageValues, void *resultValues, double return 0; } +bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo ) +{ + bool pass = true; + // Compare floats + if( memcmp( expected, actual, 4 * 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++ ) + { + if ( isnan( expected[j] ) && isnan( actual[j] ) ) + continue; + if ( IsFloatSubnormal( expected[j] ) && actual[j] == 0.0f ) + 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 46548069..1f3cfd60 100644 --- a/test_conformance/images/kernel_read_write/test_write_1D.cpp +++ b/test_conformance/images/kernel_read_write/test_write_1D.cpp @@ -19,14 +19,14 @@ #include #endif -#define MAX_ERR 0.005f - extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestMipmaps; extern cl_filter_mode gFilterModeToSkip; extern cl_mem_flags gMemFlagsToUse; extern int gtestTypesToRun; +extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo ); + const char *readwrite1DKernelSourcePattern = "__kernel void sample_kernel( __global %s4 *input, read_write image1d_t output %s)\n" "{\n" @@ -306,7 +306,7 @@ int test_write_image_1D( cl_device_id device, cl_context context, cl_command_que int numTries = 5; { char *resultPtr = (char *)resultValues; - for( size_t x = 0, i = 0; x < width_lod; x++, i++ ) + for( size_t x = 0, i = 0; i < width_lod; x++, i++ ) { char resultBuffer[ 16 ]; // Largest format would be 4 channels * 4 bytes (32 bits) each @@ -356,20 +356,14 @@ int test_write_image_1D( cl_device_id device, cl_context context, cl_command_que } else if( imageInfo->format->image_channel_data_type == CL_FLOAT ) { - // Compare floats float *expected = (float *)resultBuffer; float *actual = (float *)resultPtr; - float err = 0.f; - for( unsigned int j = 0; j < get_format_channel_count( imageInfo->format ); j++ ) - err += ( expected[ j ] != 0 ) ? fabsf( ( expected[ j ] - actual[ j ] ) / expected[ j ] ) : fabsf( expected[ j ] - actual[ j ] ); - err /= (float)get_format_channel_count( imageInfo->format ); - if( err > MAX_ERR ) + if( !validate_float_write_results( expected, actual, imageInfo ) ) { - unsigned int *e = (unsigned int *)expected; - unsigned int *a = (unsigned int *)actual; - log_error( "ERROR: Sample %ld (%ld) did not validate! (%s)\n", i, x, mem_flag_names[mem_flag_index] ); - log_error( " Error: %g\n", err ); + unsigned int *e = (unsigned int *)resultBuffer; + unsigned int *a = (unsigned int *)resultPtr; + log_error( "ERROR: Sample %ld did not validate! (%s)\n", i, mem_flag_names[ mem_flag_index ] ); log_error( " Expected: %a %a %a %a\n", expected[ 0 ], expected[ 1 ], expected[ 2 ], expected[ 3 ] ); log_error( " Expected: %08x %08x %08x %08x\n", e[ 0 ], e[ 1 ], e[ 2 ], e[ 3 ] ); log_error( " Actual: %a %a %a %a\n", actual[ 0 ], actual[ 1 ], actual[ 2 ], actual[ 3 ] ); @@ -532,7 +526,6 @@ int test_write_image_1D( cl_device_id device, cl_context context, cl_command_que return totalErrors; } - int test_write_image_1D_set( cl_device_id device, cl_context context, cl_command_queue queue, cl_image_format *format, ExplicitType inputType, MTdata d ) { char programSrc[10240]; 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 04a9b60f..c0f3e45b 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 @@ -19,13 +19,13 @@ #include #endif -#define MAX_ERR 0.005f - extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestMipmaps; extern cl_filter_mode gFilterModeToSkip; extern cl_mem_flags gMemFlagsToUse; extern int gtestTypesToRun; +extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo ); + const char *readwrite1DArrayKernelSourcePattern = "__kernel void sample_kernel( __global %s4 *input, read_write image1d_array_t output %s)\n" "{\n" @@ -375,20 +375,14 @@ int test_write_image_1D_array( cl_device_id device, cl_context context, cl_comma } else if( imageInfo->format->image_channel_data_type == CL_FLOAT ) { - // Compare floats float *expected = (float *)resultBuffer; float *actual = (float *)resultPtr; - float err = 0.f; - for( unsigned int j = 0; j < get_format_channel_count( imageInfo->format ); j++ ) - err += ( expected[ j ] != 0 ) ? fabsf( ( expected[ j ] - actual[ j ] ) / expected[ j ] ) : fabsf( expected[ j ] - actual[ j ] ); - err /= (float)get_format_channel_count( imageInfo->format ); - if( err > MAX_ERR ) + if( !validate_float_write_results( expected, actual, imageInfo ) ) { - unsigned int *e = (unsigned int *)expected; - unsigned int *a = (unsigned int *)actual; - log_error( "ERROR: Sample %ld (%ld,%ld) did not validate! (%s)\n", i, x, y, mem_flag_names[mem_flag_index] ); - log_error( " Error: %g\n", err ); + unsigned int *e = (unsigned int *)resultBuffer; + unsigned int *a = (unsigned int *)resultPtr; + log_error( "ERROR: Sample %ld (%ld,%ld) did not validate! (%s)\n", i, x, y, mem_flag_names[ mem_flag_index ] ); log_error( " Expected: %a %a %a %a\n", expected[ 0 ], expected[ 1 ], expected[ 2 ], expected[ 3 ] ); log_error( " Expected: %08x %08x %08x %08x\n", e[ 0 ], e[ 1 ], e[ 2 ], e[ 3 ] ); log_error( " Actual: %a %a %a %a\n", actual[ 0 ], actual[ 1 ], actual[ 2 ], actual[ 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 700314b8..84292fdd 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 @@ -19,16 +19,13 @@ #include #endif -#define MAX_ERR 0.005f - extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestMipmaps; extern cl_filter_mode gFilterModeToSkip; extern cl_mem_flags gMemFlagsToUse; extern int gtestTypesToRun; -extern int verify_write_results( size_t &i, int &numTries, int &totalErrors, char *&imagePtr, void *resultValues, size_t y, size_t z, - ExplicitType inputType, image_descriptor *imageInfo, bool verifyRounding ); +extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo ); // Utility function to clamp down image sizes for certain tests to avoid // using too much memory. @@ -399,20 +396,14 @@ int test_write_image_2D_array( cl_device_id device, cl_context context, cl_comma } else if( imageInfo->format->image_channel_data_type == CL_FLOAT ) { - // Compare floats float *expected = (float *)resultBuffer; float *actual = (float *)resultPtr; - float err = 0.f; - for( unsigned int j = 0; j < get_format_channel_count( imageInfo->format ); j++ ) - err += ( expected[ j ] != 0 ) ? fabsf( ( expected[ j ] - actual[ j ] ) / expected[ j ] ) : fabsf( expected[ j ] - actual[ j ] ); - err /= (float)get_format_channel_count( imageInfo->format ); - if( err > MAX_ERR ) + if( !validate_float_write_results( expected, actual, imageInfo ) ) { - unsigned int *e = (unsigned int *)expected; - unsigned int *a = (unsigned int *)actual; - log_error( "ERROR: Sample %ld (%ld,%ld) did not validate! (%s)\n", i, x, y, mem_flag_names[mem_flag_index] ); - log_error( " Error: %g\n", err ); + unsigned int *e = (unsigned int *)resultBuffer; + unsigned int *a = (unsigned int *)resultPtr; + log_error( "ERROR: Sample %ld (%ld,%ld,%ld) did not validate! (%s)\n", i, x, y, z, mem_flag_names[ mem_flag_index ] ); log_error( " Expected: %a %a %a %a\n", expected[ 0 ], expected[ 1 ], expected[ 2 ], expected[ 3 ] ); log_error( " Expected: %08x %08x %08x %08x\n", e[ 0 ], e[ 1 ], e[ 2 ], e[ 3 ] ); log_error( " Actual: %a %a %a %a\n", actual[ 0 ], actual[ 1 ], actual[ 2 ], actual[ 3 ] ); 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 f1f5bd89..14510f59 100644 --- a/test_conformance/images/kernel_read_write/test_write_3D.cpp +++ b/test_conformance/images/kernel_read_write/test_write_3D.cpp @@ -19,16 +19,13 @@ #include #endif -#define MAX_ERR 0.005f - extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestMipmaps; extern cl_filter_mode gFilterModeToSkip; extern cl_mem_flags gMemFlagsToUse; extern int gtestTypesToRun; -extern int verify_write_results( size_t &i, int &numTries, int &totalErrors, char *&imagePtr, void *resultValues, size_t y, size_t z, - ExplicitType inputType, image_descriptor *imageInfo, bool verifyRounding ); +extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo ); // Utility function to clamp down image sizes for certain tests to avoid // using too much memory. @@ -406,20 +403,14 @@ int test_write_image_3D( cl_device_id device, cl_context context, cl_command_que } else if( imageInfo->format->image_channel_data_type == CL_FLOAT ) { - // Compare floats float *expected = (float *)resultBuffer; float *actual = (float *)resultPtr; - float err = 0.f; - for( unsigned int j = 0; j < get_format_channel_count( imageInfo->format ); j++ ) - err += ( expected[ j ] != 0 ) ? fabsf( ( expected[ j ] - actual[ j ] ) / expected[ j ] ) : fabsf( expected[ j ] - actual[ j ] ); - err /= (float)get_format_channel_count( imageInfo->format ); - if( err > MAX_ERR ) + if( !validate_float_write_results( expected, actual, imageInfo ) ) { - unsigned int *e = (unsigned int *)expected; - unsigned int *a = (unsigned int *)actual; - log_error( "ERROR: Sample %ld (%ld,%ld) did not validate! (%s)\n", i, x, y, mem_flag_names[mem_flag_index] ); - log_error( " Error: %g\n", err ); + unsigned int *e = (unsigned int *)resultBuffer; + unsigned int *a = (unsigned int *)resultPtr; + log_error( "ERROR: Sample %ld (%ld,%ld,%ld) did not validate! (%s)\n", i, x, y, z, mem_flag_names[ mem_flag_index ] ); log_error( " Expected: %a %a %a %a\n", expected[ 0 ], expected[ 1 ], expected[ 2 ], expected[ 3 ] ); log_error( " Expected: %08x %08x %08x %08x\n", e[ 0 ], e[ 1 ], e[ 2 ], e[ 3 ] ); log_error( " Actual: %a %a %a %a\n", actual[ 0 ], actual[ 1 ], actual[ 2 ], actual[ 3 ] ); 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 cd689661..ffb73db1 100644 --- a/test_conformance/images/kernel_read_write/test_write_image.cpp +++ b/test_conformance/images/kernel_read_write/test_write_image.cpp @@ -19,8 +19,6 @@ #include #endif -#define MAX_ERR 0.005f - extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestImage2DFromBuffer, gTestMipmaps; extern cl_filter_mode gFilterModeToSkip; extern cl_mem_flags gMemFlagsToUse; @@ -30,7 +28,7 @@ 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 ); const char *writeKernelSourcePattern = "__kernel void sample_kernel( __global %s%s *input, write_only %s output %s)\n" @@ -422,20 +420,14 @@ int test_write_image( cl_device_id device, cl_context context, cl_command_queue } else if( imageInfo->format->image_channel_data_type == CL_FLOAT ) { - // Compare floats float *expected = (float *)resultBuffer; float *actual = (float *)resultPtr; - float err = 0.f; - for( unsigned int j = 0; j < get_format_channel_count( imageInfo->format ); j++ ) - err += ( expected[ j ] != 0 ) ? fabsf( ( expected[ j ] - actual[ j ] ) / expected[ j ] ) : fabsf( expected[ j ] - actual[ j ] ); - err /= (float)get_format_channel_count( imageInfo->format ); - if( err > MAX_ERR ) + if( !validate_float_write_results( expected, actual, imageInfo ) ) { - unsigned int *e = (unsigned int *)expected; - unsigned int *a = (unsigned int *)actual; - log_error( "ERROR: Sample %ld (%ld,%ld) did not validate! (%s)\n", i, x, y, mem_flag_names[mem_flag_index] ); - log_error( " Error: %g\n", err ); + unsigned int *e = (unsigned int *)resultBuffer; + unsigned int *a = (unsigned int *)resultPtr; + log_error( "ERROR: Sample %ld (%ld,%ld) did not validate! (%s)\n", i, x, y, mem_flag_names[ mem_flag_index ] ); log_error( " Expected: %a %a %a %a\n", expected[ 0 ], expected[ 1 ], expected[ 2 ], expected[ 3 ] ); log_error( " Expected: %08x %08x %08x %08x\n", e[ 0 ], e[ 1 ], e[ 2 ], e[ 3 ] ); log_error( " Actual: %a %a %a %a\n", actual[ 0 ], actual[ 1 ], actual[ 2 ], actual[ 3 ] );