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

Allow CL_FLOAT denorm flushing for write tests (#28) (#456)

Browse Source 
* 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
This commit is contained in:
Jim Lewis
2019-11-15 07:38:15 -06:00
committed by Kévin Petit
parent 88004077e3
commit 040321d8b9
7 changed files with 50 additions and 69 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
test_common/harness/imageHelpers.cpp
View File

@@ -1808,7 +1808,7 @@ static inline void check_for_denorms(float a[4], int *containsDenorms )
{ {
for( int i = 0; i < 4; i++ ) for( int i = 0; i < 4; i++ )
{ {
if( fabsf(a[i]) < FLT_MIN ) if( IsFloatSubnormal( a[i] ) )
a[i] = copysignf( 0.0f, 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++ ) for( int i = 0; i < 4; i++ )
{ {
if( fabs(a[i]) < FLT_MIN ) if( IsFloatSubnormal( a[i] ) )
{ {
*containsDenorms = 1; *containsDenorms = 1;
break; break;

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

@@ -1151,6 +1151,26 @@ int validate_image_2D_sRGB_results(void *imageValues, void *resultValues, double
return 0; 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, int test_read_image_2D( cl_context context, cl_command_queue queue, cl_kernel kernel,
image_descriptor *imageInfo, image_sampler_data *imageSampler, image_descriptor *imageInfo, image_sampler_data *imageSampler,
bool useFloatCoords, ExplicitType outputType, MTdata d ) bool useFloatCoords, ExplicitType outputType, MTdata d )

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

@@ -19,14 +19,14 @@
#include <sys/mman.h> #include <sys/mman.h>
#endif #endif
#define MAX_ERR 0.005f
extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestMipmaps; extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestMipmaps;
extern cl_filter_mode gFilterModeToSkip; extern cl_filter_mode gFilterModeToSkip;
extern cl_mem_flags gMemFlagsToUse; extern cl_mem_flags gMemFlagsToUse;
extern int gtestTypesToRun; extern int gtestTypesToRun;
extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo );
const char *readwrite1DKernelSourcePattern = const char *readwrite1DKernelSourcePattern =
"__kernel void sample_kernel( __global %s4 *input, read_write image1d_t output %s)\n" "__kernel void sample_kernel( __global %s4 *input, read_write image1d_t output %s)\n"
"{\n" "{\n"
@@ -306,7 +306,7 @@ int test_write_image_1D( cl_device_id device, cl_context context, cl_command_que
int numTries = 5; int numTries = 5;
{ {
char *resultPtr = (char *)resultValues; 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 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 ) else if( imageInfo->format->image_channel_data_type == CL_FLOAT )
{ {
// Compare floats
float *expected = (float *)resultBuffer; float *expected = (float *)resultBuffer;
float *actual = (float *)resultPtr; 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( !validate_float_write_results( expected, actual, imageInfo ) )
if( err > MAX_ERR )
{ {
unsigned int *e = (unsigned int *)expected; unsigned int *e = (unsigned int *)resultBuffer;
unsigned int *a = (unsigned int *)actual; unsigned int *a = (unsigned int *)resultPtr;
log_error( "ERROR: Sample %ld (%ld) did not validate! (%s)\n", i, x, mem_flag_names[mem_flag_index] ); log_error( "ERROR: Sample %ld did not validate! (%s)\n", i, mem_flag_names[ mem_flag_index ] );
log_error( " Error: %g\n", err );
log_error( " Expected: %a %a %a %a\n", expected[ 0 ], expected[ 1 ], expected[ 2 ], expected[ 3 ] ); 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( " 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 ] ); 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; 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 ) 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]; char programSrc[10240];

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

@@ -19,13 +19,13 @@
#include <sys/mman.h> #include <sys/mman.h>
#endif #endif
#define MAX_ERR 0.005f
extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestMipmaps; extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestMipmaps;
extern cl_filter_mode gFilterModeToSkip; extern cl_filter_mode gFilterModeToSkip;
extern cl_mem_flags gMemFlagsToUse; extern cl_mem_flags gMemFlagsToUse;
extern int gtestTypesToRun; extern int gtestTypesToRun;
extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo );
const char *readwrite1DArrayKernelSourcePattern = const char *readwrite1DArrayKernelSourcePattern =
"__kernel void sample_kernel( __global %s4 *input, read_write image1d_array_t output %s)\n" "__kernel void sample_kernel( __global %s4 *input, read_write image1d_array_t output %s)\n"
"{\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 ) else if( imageInfo->format->image_channel_data_type == CL_FLOAT )
{ {
// Compare floats
float *expected = (float *)resultBuffer; float *expected = (float *)resultBuffer;
float *actual = (float *)resultPtr; 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( !validate_float_write_results( expected, actual, imageInfo ) )
if( err > MAX_ERR )
{ {
unsigned int *e = (unsigned int *)expected; unsigned int *e = (unsigned int *)resultBuffer;
unsigned int *a = (unsigned int *)actual; 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( "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 );
log_error( " Expected: %a %a %a %a\n", expected[ 0 ], expected[ 1 ], expected[ 2 ], expected[ 3 ] ); 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( " 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 ] ); log_error( " Actual: %a %a %a %a\n", actual[ 0 ], actual[ 1 ], actual[ 2 ], actual[ 3 ] );

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

@@ -19,16 +19,13 @@
#include <sys/mman.h> #include <sys/mman.h>
#endif #endif
#define MAX_ERR 0.005f
extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestMipmaps; extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestMipmaps;
extern cl_filter_mode gFilterModeToSkip; extern cl_filter_mode gFilterModeToSkip;
extern cl_mem_flags gMemFlagsToUse; extern cl_mem_flags gMemFlagsToUse;
extern int gtestTypesToRun; 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, extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo );
ExplicitType inputType, image_descriptor *imageInfo, bool verifyRounding );
// Utility function to clamp down image sizes for certain tests to avoid // Utility function to clamp down image sizes for certain tests to avoid
// using too much memory. // 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 ) else if( imageInfo->format->image_channel_data_type == CL_FLOAT )
{ {
// Compare floats
float *expected = (float *)resultBuffer; float *expected = (float *)resultBuffer;
float *actual = (float *)resultPtr; 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( !validate_float_write_results( expected, actual, imageInfo ) )
if( err > MAX_ERR )
{ {
unsigned int *e = (unsigned int *)expected; unsigned int *e = (unsigned int *)resultBuffer;
unsigned int *a = (unsigned int *)actual; 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( "ERROR: Sample %ld (%ld,%ld,%ld) did not validate! (%s)\n", i, x, y, z, mem_flag_names[ mem_flag_index ] );
log_error( " Error: %g\n", err );
log_error( " Expected: %a %a %a %a\n", expected[ 0 ], expected[ 1 ], expected[ 2 ], expected[ 3 ] ); 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( " 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 ] ); log_error( " Actual: %a %a %a %a\n", actual[ 0 ], actual[ 1 ], actual[ 2 ], actual[ 3 ] );

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

@@ -19,16 +19,13 @@
#include <sys/mman.h> #include <sys/mman.h>
#endif #endif
#define MAX_ERR 0.005f
extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestMipmaps; extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestMipmaps;
extern cl_filter_mode gFilterModeToSkip; extern cl_filter_mode gFilterModeToSkip;
extern cl_mem_flags gMemFlagsToUse; extern cl_mem_flags gMemFlagsToUse;
extern int gtestTypesToRun; 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, extern bool validate_float_write_results( float *expected, float *actual, image_descriptor *imageInfo );
ExplicitType inputType, image_descriptor *imageInfo, bool verifyRounding );
// Utility function to clamp down image sizes for certain tests to avoid // Utility function to clamp down image sizes for certain tests to avoid
// using too much memory. // 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 ) else if( imageInfo->format->image_channel_data_type == CL_FLOAT )
{ {
// Compare floats
float *expected = (float *)resultBuffer; float *expected = (float *)resultBuffer;
float *actual = (float *)resultPtr; 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( !validate_float_write_results( expected, actual, imageInfo ) )
if( err > MAX_ERR )
{ {
unsigned int *e = (unsigned int *)expected; unsigned int *e = (unsigned int *)resultBuffer;
unsigned int *a = (unsigned int *)actual; 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( "ERROR: Sample %ld (%ld,%ld,%ld) did not validate! (%s)\n", i, x, y, z, mem_flag_names[ mem_flag_index ] );
log_error( " Error: %g\n", err );
log_error( " Expected: %a %a %a %a\n", expected[ 0 ], expected[ 1 ], expected[ 2 ], expected[ 3 ] ); 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( " 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 ] ); log_error( " Actual: %a %a %a %a\n", actual[ 0 ], actual[ 1 ], actual[ 2 ], actual[ 3 ] );

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

@@ -19,8 +19,6 @@
#include <sys/mman.h> #include <sys/mman.h>
#endif #endif
#define MAX_ERR 0.005f
extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestImage2DFromBuffer, gTestMipmaps; extern bool gDebugTrace, gDisableOffsets, gTestSmallImages, gEnablePitch, gTestMaxImages, gTestRounding, gTestImage2DFromBuffer, gTestMipmaps;
extern cl_filter_mode gFilterModeToSkip; extern cl_filter_mode gFilterModeToSkip;
extern cl_mem_flags gMemFlagsToUse; 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_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_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 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 = const char *writeKernelSourcePattern =
"__kernel void sample_kernel( __global %s%s *input, write_only %s output %s)\n" "__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 ) else if( imageInfo->format->image_channel_data_type == CL_FLOAT )
{ {
// Compare floats
float *expected = (float *)resultBuffer; float *expected = (float *)resultBuffer;
float *actual = (float *)resultPtr; 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( !validate_float_write_results( expected, actual, imageInfo ) )
if( err > MAX_ERR )
{ {
unsigned int *e = (unsigned int *)expected; unsigned int *e = (unsigned int *)resultBuffer;
unsigned int *a = (unsigned int *)actual; 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( "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 );
log_error( " Expected: %a %a %a %a\n", expected[ 0 ], expected[ 1 ], expected[ 2 ], expected[ 3 ] ); 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( " 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 ] ); log_error( " Actual: %a %a %a %a\n", actual[ 0 ], actual[ 1 ], actual[ 2 ], actual[ 3 ] );