From 944b0a817823617d7b65fc01b03d27bf8de91433 Mon Sep 17 00:00:00 2001 From: James Price Date: Wed, 27 May 2020 14:13:11 -0400 Subject: [PATCH] Enable -Werror for GCC/Clang builds (#786) * Enable -Werror for GCC/Clang builds Fixes many of the errors this produces, and disables a handful that didn't have solutions that were obvious (to me). * Check for `-W*` flags empirically * Remove cl_APPLE_fp64_basic_ops support * Undo NAN conversion fix * Add comments to warning override flags * Remove unneeded STRINGIFY definition * Fix tautological compare issue in basic * Use ABS_ERROR macro in image tests * Use fabs for ABS_ERROR macro * Move ABS_ERROR definition to common header --- CMakeLists.txt | 20 ++- test_common/harness/imageHelpers.cpp | 3 +- test_conformance/api/test_null_buffer_arg.cpp | 2 +- .../basic/test_work_item_functions.cpp | 6 +- .../buffers/test_buffer_migrate.cpp | 12 +- .../buffers/test_image_migrate.cpp | 12 +- .../compiler/test_image_macro.cpp | 2 +- test_conformance/computeinfo/main.cpp | 2 +- test_conformance/half/cl_utils.h | 3 - .../clCopyImage/test_copy_2D_2D_array.cpp | 8 +- .../images/clCopyImage/test_copy_2D_3D.cpp | 8 +- .../images/kernel_image_methods/test_1D.cpp | 2 +- .../kernel_image_methods/test_1D_array.cpp | 2 +- .../images/kernel_image_methods/test_2D.cpp | 2 +- .../kernel_image_methods/test_2D_array.cpp | 2 +- .../kernel_image_methods/test_loops.cpp | 5 +- .../images/kernel_read_write/test_common.h | 2 + .../kernel_read_write/test_iterations.cpp | 86 ++++++------ .../images/kernel_read_write/test_read_1D.cpp | 88 ++++++++----- .../kernel_read_write/test_read_1D_array.cpp | 88 ++++++++----- .../kernel_read_write/test_read_2D_array.cpp | 122 ++++++++++++------ .../images/kernel_read_write/test_read_3D.cpp | 114 ++++++++++------ test_conformance/math_brute_force/main.cpp | 121 ++--------------- .../math_brute_force/reference_math.cpp | 33 ----- test_conformance/spir/datagen.h | 14 +- test_conformance/spir/main.cpp | 4 +- test_conformance/spir/run_build_test.cpp | 2 +- 27 files changed, 398 insertions(+), 367 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index f811b35d..799460da 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -124,6 +124,7 @@ endif (GL_IS_SUPPORTED AND CLConform_GL_LIBRARIES_DIR) include(CheckFunctionExists) include(CheckIncludeFiles) +include(CheckCXXCompilerFlag) if(CMAKE_SYSTEM_PROCESSOR MATCHES "^(arm.*|ARM.*)") set(CLConform_TARGET_ARCH ARM) @@ -139,9 +140,24 @@ if(NOT DEFINED CLConform_TARGET_ARCH) message (FATAL_ERROR "Target architecture not recognised. Exiting.") endif() +macro(add_cxx_flag_if_supported flag) + string(REGEX REPLACE "[-=+]" "" FLAG_NO_SIGNS ${flag}) + check_cxx_compiler_flag(${flag} COMPILER_SUPPORTS_${FLAG_NO_SIGNS}) + if(COMPILER_SUPPORTS_${FLAG_NO_SIGNS}) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${flag}") + endif() +endmacro(add_cxx_flag_if_supported) + if(CMAKE_COMPILER_IS_GNUCC OR "${CMAKE_CXX_COMPILER_ID}" MATCHES "(Apple)?Clang") - set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -Wno-narrowing") - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-narrowing") + add_cxx_flag_if_supported(-Wno-narrowing) + add_cxx_flag_if_supported(-Wno-format) + add_cxx_flag_if_supported(-Werror) + add_cxx_flag_if_supported(-Wno-error=cpp) # Allow #warning directive + add_cxx_flag_if_supported(-Wno-error=absolute-value) # Issue 783 + add_cxx_flag_if_supported(-Wno-error=unknown-pragmas) # Issue #785 + add_cxx_flag_if_supported(-Wno-error=asm-operand-widths) # Issue #784 + add_cxx_flag_if_supported(-Wno-error=overflow) # Fixed by #699 + # -msse -mfpmath=sse to force gcc to use sse for float math, # avoiding excess precision problems that cause tests like int2float # to falsely fail. -ffloat-store also works, but WG suggested diff --git a/test_common/harness/imageHelpers.cpp b/test_common/harness/imageHelpers.cpp index 75ce37ed..f36c153e 100644 --- a/test_common/harness/imageHelpers.cpp +++ b/test_common/harness/imageHelpers.cpp @@ -899,8 +899,7 @@ int get_format_min_int( cl_image_format *format ) case CL_UNORM_INT_101010: return 0; - case CL_HALF_FLOAT: - return -1<<10; + case CL_HALF_FLOAT: return -(1 << 10); #ifdef CL_SFIXED14_APPLE case CL_SFIXED14_APPLE: diff --git a/test_conformance/api/test_null_buffer_arg.cpp b/test_conformance/api/test_null_buffer_arg.cpp index 4fa046b6..ba43f183 100644 --- a/test_conformance/api/test_null_buffer_arg.cpp +++ b/test_conformance/api/test_null_buffer_arg.cpp @@ -62,7 +62,7 @@ static int test_setargs_and_execution(cl_command_queue queue, cl_kernel kernel, unsigned int i; cl_int status; - char *typestr; + const char *typestr; if (type == NON_NULL_PATH) { status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &test_buf); diff --git a/test_conformance/basic/test_work_item_functions.cpp b/test_conformance/basic/test_work_item_functions.cpp index 890f3c2d..d95915cf 100644 --- a/test_conformance/basic/test_work_item_functions.cpp +++ b/test_conformance/basic/test_work_item_functions.cpp @@ -130,7 +130,7 @@ int test_work_item_functions(cl_device_id deviceID, cl_context context, cl_comma free_mtdata(d); return -1; } - if( testData[ q ].globalID[ j ] < 0 || testData[ q ].globalID[ j ] >= (cl_uint)threads[ j ] ) + if (testData[q].globalID[j] >= (cl_uint)threads[j]) { log_error( "ERROR: get_global_id(%d) did not return proper value for %d dimensions (max %d, got %d)\n", (int)j, (int)dim, (int)threads[ j ], (int)testData[ q ].globalID[ j ] ); @@ -144,7 +144,7 @@ int test_work_item_functions(cl_device_id deviceID, cl_context context, cl_comma free_mtdata(d); return -1; } - if( testData[ q ].localID[ j ] < 0 && testData[ q ].localID[ j ] >= (cl_uint)localThreads[ j ] ) + if (testData[q].localID[j] >= (cl_uint)localThreads[j]) { log_error( "ERROR: get_local_id(%d) did not return proper value for %d dimensions (max %d, got %d)\n", (int)j, (int)dim, (int)localThreads[ j ], (int)testData[ q ].localID[ j ] ); @@ -159,7 +159,7 @@ int test_work_item_functions(cl_device_id deviceID, cl_context context, cl_comma free_mtdata(d); return -1; } - if( testData[ q ].groupID[ j ] < 0 || testData[ q ].groupID[ j ] >= (cl_uint)groupCount ) + if (testData[q].groupID[j] >= (cl_uint)groupCount) { log_error( "ERROR: get_group_id(%d) did not return proper value for %d dimensions (max %d, got %d)\n", (int)j, (int)dim, (int)groupCount, (int)testData[ q ].groupID[ j ] ); diff --git a/test_conformance/buffers/test_buffer_migrate.cpp b/test_conformance/buffers/test_buffer_migrate.cpp index 1ae3f1e1..a5b6f26a 100644 --- a/test_conformance/buffers/test_buffer_migrate.cpp +++ b/test_conformance/buffers/test_buffer_migrate.cpp @@ -69,6 +69,7 @@ static cl_int migrateMemObject(enum migrations migrate, cl_command_queue *queues // Choose a random set of flags flags[i] = (cl_mem_migration_flags)(genrand_int32(d) & (CL_MIGRATE_MEM_OBJECT_HOST | CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED));; break; + default: log_error("Unhandled migration type: %d\n", migrate); return -1; } if ((err = clEnqueueMigrateMemObjects(queues[j], 1, (const cl_mem *)(&mem_objects[i]), flags[i], 0, NULL, NULL)) != CL_SUCCESS) { print_error(err, "Failed migrating memory object."); @@ -218,10 +219,13 @@ int test_buffer_migrate(cl_device_id deviceID, cl_context context, cl_command_qu } // Build the kernel program. - if (err = create_single_kernel_helper(ctx, &program, &kernel, 1, &buffer_migrate_kernel_code, "test_buffer_migrate")) { - print_error(err, "Failed creating kernel."); - failed = 1; - goto cleanup; + if ((err = create_single_kernel_helper(ctx, &program, &kernel, 1, + &buffer_migrate_kernel_code, + "test_buffer_migrate"))) + { + print_error(err, "Failed creating kernel."); + failed = 1; + goto cleanup; } num_devices_limited = num_devices; diff --git a/test_conformance/buffers/test_image_migrate.cpp b/test_conformance/buffers/test_image_migrate.cpp index f381f374..31bb0a21 100644 --- a/test_conformance/buffers/test_image_migrate.cpp +++ b/test_conformance/buffers/test_image_migrate.cpp @@ -80,6 +80,7 @@ static cl_int migrateMemObject(enum migrations migrate, cl_command_queue *queues // Choose a random set of flags flags[i] = (cl_mem_migration_flags)(genrand_int32(d) & (CL_MIGRATE_MEM_OBJECT_HOST | CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED)); break; + default: log_error("Unhandled migration type: %d\n", migrate); return -1; } if ((err = clEnqueueMigrateMemObjects(queues[j], 1, (const cl_mem *)(&mem_objects[i]), flags[i], 0, NULL, NULL)) != CL_SUCCESS) { @@ -251,10 +252,13 @@ int test_image_migrate(cl_device_id deviceID, cl_context context, cl_command_que } // Build the kernel program. - if (err = create_single_kernel_helper(ctx, &program, &kernel, 1, &image_migrate_kernel_code, "test_image_migrate")) { - print_error(err, "Failed creating kernel."); - failed = 1; - goto cleanup; + if ((err = create_single_kernel_helper(ctx, &program, &kernel, 1, + &image_migrate_kernel_code, + "test_image_migrate"))) + { + print_error(err, "Failed creating kernel."); + failed = 1; + goto cleanup; } // Create sampler. diff --git a/test_conformance/compiler/test_image_macro.cpp b/test_conformance/compiler/test_image_macro.cpp index f85f092f..8a5c6707 100644 --- a/test_conformance/compiler/test_image_macro.cpp +++ b/test_conformance/compiler/test_image_macro.cpp @@ -57,7 +57,7 @@ int test_image_macro(cl_device_id deviceID, cl_context context, cl_command_queue return status; } - if( (image_support == CL_TRUE) ) + if (image_support == CL_TRUE) { status = create_single_kernel_helper_create_program(context, &program, 1, (const char**)&image_supported_source); diff --git a/test_conformance/computeinfo/main.cpp b/test_conformance/computeinfo/main.cpp index 5900ffda..0bc04a9c 100644 --- a/test_conformance/computeinfo/main.cpp +++ b/test_conformance/computeinfo/main.cpp @@ -96,7 +96,7 @@ typedef struct _extensions extensions_t; // first is greater). int vercmp(version_t a, version_t b) { - if (a.major < b.major || a.major == b.major && a.minor < b.minor) + if (a.major < b.major || (a.major == b.major && a.minor < b.minor)) { return -1; } diff --git a/test_conformance/half/cl_utils.h b/test_conformance/half/cl_utils.h index 814cb902..82a63115 100644 --- a/test_conformance/half/cl_utils.h +++ b/test_conformance/half/cl_utils.h @@ -93,9 +93,6 @@ void ReleaseCL( void ); int RunKernel( cl_device_id device, cl_kernel kernel, void *inBuf, void *outBuf, uint32_t blockCount , int extraArg); cl_program MakeProgram( cl_device_id device, const char *source[], int count ); -#define STRING( _x ) STRINGIFY( _x ) -#define STRINGIFY(x) #x - static inline float as_float(cl_uint u) { union { cl_uint u; float f; }v; v.u = u; return v.f; } static inline double as_double(cl_ulong u) { union { cl_ulong u; double d; }v; v.u = u; return v.d; } diff --git a/test_conformance/images/clCopyImage/test_copy_2D_2D_array.cpp b/test_conformance/images/clCopyImage/test_copy_2D_2D_array.cpp index 75cf1af4..f784230f 100644 --- a/test_conformance/images/clCopyImage/test_copy_2D_2D_array.cpp +++ b/test_conformance/images/clCopyImage/test_copy_2D_2D_array.cpp @@ -166,7 +166,8 @@ int test_copy_image_size_2D_2D_array( cl_context context, cl_command_queue queue sourcePos[ 0 ] = ( srcImageInfo->width > regionSize[ 0 ] ) ? (size_t)random_in_range( 0, (int)( srcImageInfo->width - regionSize[ 0 ] - 1 ), d ) : 0; sourcePos[ 1 ] = ( srcImageInfo->height > regionSize[ 1 ] ) ? (size_t)random_in_range( 0, (int)( srcImageInfo->height - regionSize[ 1 ] - 1 ), d ) : 0; sourcePos[ 2 ] = ( srcImageInfo->arraySize > 0 ) ? (size_t)random_in_range( 0, (int)( srcImageInfo->arraySize - 1 ), d ) : gTestMipmaps ? twoImage_lod : 0; - if ( gTestMipmaps ) + if (gTestMipmaps) + { if( srcImageInfo->arraySize > 0 ) { sourcePos[ 0 ] = ( threeImage_width_lod > regionSize[ 0 ] ) ? (size_t)random_in_range( 0, (int)( threeImage_width_lod - regionSize[ 0 ] - 1 ), d ) : 0; @@ -179,11 +180,13 @@ int test_copy_image_size_2D_2D_array( cl_context context, cl_command_queue queue sourcePos[ 1 ] = ( twoImage_height_lod > regionSize[ 1 ] ) ? (size_t)random_in_range( 0, (int)( twoImage_height_lod - regionSize[ 1 ] - 1 ), d ) : 0; } + } destPos[ 0 ] = ( dstImageInfo->width > regionSize[ 0 ] ) ? (size_t)random_in_range( 0, (int)( dstImageInfo->width - regionSize[ 0 ] - 1 ), d ) : 0; destPos[ 1 ] = ( dstImageInfo->height > regionSize[ 1 ] ) ? (size_t)random_in_range( 0, (int)( dstImageInfo->height - regionSize[ 1 ] - 1 ), d ) : 0; destPos[ 2 ] = ( dstImageInfo->arraySize > 0 ) ? (size_t)random_in_range( 0, (int)( dstImageInfo->arraySize - 1 ), d ) : gTestMipmaps ? twoImage_lod : 0; - if ( gTestMipmaps ) + if (gTestMipmaps) + { if( dstImageInfo->arraySize > 0 ) { destPos[ 0 ] = ( threeImage_width_lod > regionSize[ 0 ] ) ? (size_t)random_in_range( 0, (int)( threeImage_width_lod - regionSize[ 0 ] - 1 ), d ) : 0; @@ -196,6 +199,7 @@ int test_copy_image_size_2D_2D_array( cl_context context, cl_command_queue queue destPos[ 1 ] = ( twoImage_height_lod > regionSize[ 1 ] ) ? (size_t)random_in_range( 0, (int)( twoImage_height_lod - regionSize[ 1 ] - 1 ), d ) : 0; } + } // Go for it! retCode = test_copy_image_generic( context, queue, srcImageInfo, dstImageInfo, sourcePos, destPos, regionSize, d ); diff --git a/test_conformance/images/clCopyImage/test_copy_2D_3D.cpp b/test_conformance/images/clCopyImage/test_copy_2D_3D.cpp index 834ea4f4..d341455e 100644 --- a/test_conformance/images/clCopyImage/test_copy_2D_3D.cpp +++ b/test_conformance/images/clCopyImage/test_copy_2D_3D.cpp @@ -167,7 +167,8 @@ int test_copy_image_size_2D_3D( cl_context context, cl_command_queue queue, imag sourcePos[ 1 ] = ( srcImageInfo->height > regionSize[ 1 ] ) ? (size_t)random_in_range( 0, (int)( srcImageInfo->height - regionSize[ 1 ] - 1 ), d ) : 0; sourcePos[ 2 ] = ( srcImageInfo->depth > 0 ) ? (size_t)random_in_range( 0, (int)( srcImageInfo->depth - 1 ), d ) : 0; - if ( gTestMipmaps ) + if (gTestMipmaps) + { if( srcImageInfo->depth > 0 ) { sourcePos[ 0 ] = ( threeImage_width_lod > regionSize[ 0 ] ) ? (size_t)random_in_range( 0, (int)( threeImage_width_lod - regionSize[ 0 ] - 1 ), d ) : 0; @@ -181,12 +182,14 @@ int test_copy_image_size_2D_3D( cl_context context, cl_command_queue queue, imag sourcePos[ 1 ] = ( twoImage_height_lod > regionSize[ 1 ] ) ? (size_t)random_in_range( 0, (int)( twoImage_height_lod - regionSize[ 1 ] - 1 ), d ) : 0; } + } destPos[ 0 ] = ( dstImageInfo->width > regionSize[ 0 ] ) ? (size_t)random_in_range( 0, (int)( dstImageInfo->width - regionSize[ 0 ] - 1 ), d ) : 0; destPos[ 1 ] = ( dstImageInfo->height > regionSize[ 1 ] ) ? (size_t)random_in_range( 0, (int)( dstImageInfo->height - regionSize[ 1 ] - 1 ), d ) : 0; destPos[ 2 ] = ( dstImageInfo->depth > 0 ) ? (size_t)random_in_range( 0, (int)( dstImageInfo->depth - 1 ), d ) : 0; - if ( gTestMipmaps ) + if (gTestMipmaps) + { if( dstImageInfo->depth > 0 ) { destPos[ 0 ] = ( threeImage_width_lod > regionSize[ 0 ] ) ? (size_t)random_in_range( 0, (int)( threeImage_width_lod - regionSize[ 0 ] - 1 ), d ) : 0; @@ -200,6 +203,7 @@ int test_copy_image_size_2D_3D( cl_context context, cl_command_queue queue, imag destPos[ 1 ] = ( twoImage_height_lod > regionSize[ 1 ] ) ? (size_t)random_in_range( 0, (int)( twoImage_height_lod - regionSize[ 1 ] - 1 ), d ) : 0; } + } // Go for it! retCode = test_copy_image_generic( context, queue, srcImageInfo, dstImageInfo, sourcePos, destPos, regionSize, d ); diff --git a/test_conformance/images/kernel_image_methods/test_1D.cpp b/test_conformance/images/kernel_image_methods/test_1D.cpp index 07e9cf42..757a4a00 100644 --- a/test_conformance/images/kernel_image_methods/test_1D.cpp +++ b/test_conformance/images/kernel_image_methods/test_1D.cpp @@ -20,7 +20,7 @@ extern bool gDebugTrace, gTestSmallImages, gTestMaxImages, gDeviceLt20; -typedef struct image_kernel_data +struct image_kernel_data { cl_int width; cl_int channelType; diff --git a/test_conformance/images/kernel_image_methods/test_1D_array.cpp b/test_conformance/images/kernel_image_methods/test_1D_array.cpp index 83594a2b..f5e778b6 100644 --- a/test_conformance/images/kernel_image_methods/test_1D_array.cpp +++ b/test_conformance/images/kernel_image_methods/test_1D_array.cpp @@ -20,7 +20,7 @@ extern bool gDebugTrace, gTestSmallImages, gTestMaxImages, gDeviceLt20; -typedef struct image_kernel_data +struct image_kernel_data { cl_int width; cl_int arraySize; diff --git a/test_conformance/images/kernel_image_methods/test_2D.cpp b/test_conformance/images/kernel_image_methods/test_2D.cpp index ab9138d5..64b9f265 100644 --- a/test_conformance/images/kernel_image_methods/test_2D.cpp +++ b/test_conformance/images/kernel_image_methods/test_2D.cpp @@ -20,7 +20,7 @@ extern bool gDebugTrace, gTestSmallImages, gTestMaxImages, gDeviceLt20; -typedef struct image_kernel_data +struct image_kernel_data { cl_int width; cl_int height; diff --git a/test_conformance/images/kernel_image_methods/test_2D_array.cpp b/test_conformance/images/kernel_image_methods/test_2D_array.cpp index f621cd6d..85b8a7a3 100644 --- a/test_conformance/images/kernel_image_methods/test_2D_array.cpp +++ b/test_conformance/images/kernel_image_methods/test_2D_array.cpp @@ -20,7 +20,7 @@ extern bool gDebugTrace, gTestSmallImages, gTestMaxImages, gDeviceLt20; -typedef struct image_kernel_data +struct image_kernel_data { cl_int width; cl_int height; diff --git a/test_conformance/images/kernel_image_methods/test_loops.cpp b/test_conformance/images/kernel_image_methods/test_loops.cpp index 5465dcb6..3b56d3ef 100644 --- a/test_conformance/images/kernel_image_methods/test_loops.cpp +++ b/test_conformance/images/kernel_image_methods/test_loops.cpp @@ -111,7 +111,8 @@ int test_image_set( cl_device_id device, cl_context context, cl_command_queue qu gDeviceLt20 = true; } - if (version_check = (version < Version(1,2))) { + if ((version_check = (version < Version(1, 2)))) + { switch (imageType) { case CL_MEM_OBJECT_IMAGE1D: test_missing_feature(version_check, "image_1D"); @@ -120,7 +121,7 @@ int test_image_set( cl_device_id device, cl_context context, cl_command_queue qu case CL_MEM_OBJECT_IMAGE2D_ARRAY: test_missing_feature(version_check, "image_2D_array"); } - } + } int ret = 0; ret += test_image_type( device, context, queue, imageType, CL_MEM_READ_ONLY ); diff --git a/test_conformance/images/kernel_read_write/test_common.h b/test_conformance/images/kernel_read_write/test_common.h index 9cacb1e6..1a1a8a1a 100644 --- a/test_conformance/images/kernel_read_write/test_common.h +++ b/test_conformance/images/kernel_read_write/test_common.h @@ -1,5 +1,7 @@ #include "../testBase.h" +#define ABS_ERROR(result, expected) (fabs(expected - result)) + extern cl_sampler create_sampler(cl_context context, image_sampler_data *sdata, bool test_mipmaps, cl_int *error); diff --git a/test_conformance/images/kernel_read_write/test_iterations.cpp b/test_conformance/images/kernel_read_write/test_iterations.cpp index 050a805d..0b7d4243 100644 --- a/test_conformance/images/kernel_read_write/test_iterations.cpp +++ b/test_conformance/images/kernel_read_write/test_iterations.cpp @@ -85,8 +85,6 @@ static const char *lodOffsetSource = static const char *offsetSource = " int offset = tidY*get_image_width(input) + tidX;\n"; -#define ABS_ERROR( result, expected ) ( fabsf( (float)expected - (float)result ) ) - extern void read_image_pixel_float( void *imageData, image_descriptor *imageInfo, int x, int y, int z, float *outData ); template int determine_validation_error( void *imagePtr, image_descriptor *imageInfo, image_sampler_data *imageSampler, @@ -453,7 +451,7 @@ int validate_image_2D_depth_results(void *imageValues, void *resultValues, doubl xOffsetValues[ j ], yOffsetValues[ j ], 0.0f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, &containsDenormals ); - float err1 = fabsf( resultPtr[0] - expected[0] ); + float err1 = ABS_ERROR(resultPtr[0], expected[0]); // Clamp to the minimum absolute error for the format if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN ); @@ -472,7 +470,7 @@ int validate_image_2D_depth_results(void *imageValues, void *resultValues, doubl xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, NULL ); - err1 = fabsf( resultPtr[0] - expected[0] ); + err1 = ABS_ERROR(resultPtr[0], expected[0]); } } @@ -505,7 +503,7 @@ int validate_image_2D_depth_results(void *imageValues, void *resultValues, doubl xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, &containsDenormals ); - float err1 = fabsf( resultPtr[0] - expected[0] ); + float err1 = ABS_ERROR(resultPtr[0], expected[0]); float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN ); @@ -520,7 +518,7 @@ int validate_image_2D_depth_results(void *imageValues, void *resultValues, doubl xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, NULL ); - err1 = fabsf( resultPtr[0] - expected[0] ); + err1 = ABS_ERROR(resultPtr[0], expected[0]); } } if( ! (err1 <= maxErr1) ) @@ -611,10 +609,10 @@ int validate_image_2D_results(void *imageValues, void *resultValues, double form xOffsetValues[ j ], yOffsetValues[ j ], 0.0f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, &containsDenormals ); - float err1 = fabsf( resultPtr[0] - expected[0] ); - float err2 = fabsf( resultPtr[1] - expected[1] ); - float err3 = fabsf( resultPtr[2] - expected[2] ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = ABS_ERROR(resultPtr[0], expected[0]); + float err2 = ABS_ERROR(resultPtr[1], expected[1]); + float err3 = ABS_ERROR(resultPtr[2], expected[2]); + float err4 = ABS_ERROR(resultPtr[3], expected[3]); // Clamp to the minimum absolute error for the format if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 0.0f; } @@ -648,10 +646,10 @@ int validate_image_2D_results(void *imageValues, void *resultValues, double form xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, NULL ); - err1 = fabsf( resultPtr[0] - expected[0] ); - err2 = fabsf( resultPtr[1] - expected[1] ); - err3 = fabsf( resultPtr[2] - expected[2] ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(resultPtr[0], expected[0]); + err2 = ABS_ERROR(resultPtr[1], expected[1]); + err3 = ABS_ERROR(resultPtr[2], expected[2]); + err4 = ABS_ERROR(resultPtr[3], expected[3]); } } @@ -689,10 +687,10 @@ int validate_image_2D_results(void *imageValues, void *resultValues, double form xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, &containsDenormals ); - float err1 = fabsf( resultPtr[0] - expected[0] ); - float err2 = fabsf( resultPtr[1] - expected[1] ); - float err3 = fabsf( resultPtr[2] - expected[2] ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = ABS_ERROR(resultPtr[0], expected[0]); + float err2 = ABS_ERROR(resultPtr[1], expected[1]); + float err3 = ABS_ERROR(resultPtr[2], expected[2]); + float err4 = ABS_ERROR(resultPtr[3], expected[3]); float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN ); float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN ); float maxErr3 = MAX( maxErr * maxPixel.p[2], FLT_MIN ); @@ -719,10 +717,10 @@ int validate_image_2D_results(void *imageValues, void *resultValues, double form xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, NULL ); - err1 = fabsf( resultPtr[0] - expected[0] ); - err2 = fabsf( resultPtr[1] - expected[1] ); - err3 = fabsf( resultPtr[2] - expected[2] ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(resultPtr[0], expected[0]); + err2 = ABS_ERROR(resultPtr[1], expected[1]); + err3 = ABS_ERROR(resultPtr[2], expected[2]); + err4 = ABS_ERROR(resultPtr[3], expected[3]); } } if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) || @@ -1008,10 +1006,13 @@ int validate_image_2D_sRGB_results(void *imageValues, void *resultValues, double maxPixel = sample_image_pixel_float_offset( imageValues, imageInfo, xOffsetValues[ j ], yOffsetValues[ j ], 0.0f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, &containsDenormals ); - float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + float err2 = ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + float err3 = ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + float err4 = ABS_ERROR(resultPtr[3], expected[3]); float maxErr = 0.5; // Check if the result matches. @@ -1034,10 +1035,13 @@ int validate_image_2D_sRGB_results(void *imageValues, void *resultValues, double xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, NULL ); - err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + err2 = ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + err3 = ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + err4 = ABS_ERROR(resultPtr[3], expected[3]); } } @@ -1075,10 +1079,13 @@ int validate_image_2D_sRGB_results(void *imageValues, void *resultValues, double xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, &containsDenormals ); - float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + float err2 = ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + float err3 = ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + float err4 = ABS_ERROR(resultPtr[3], expected[3]); float maxErr = 0.6; if( ! (err1 <= maxErr) || ! (err2 <= maxErr) || @@ -1099,10 +1106,13 @@ int validate_image_2D_sRGB_results(void *imageValues, void *resultValues, double xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, NULL ); - err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + err2 = ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + err3 = ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + err4 = ABS_ERROR(resultPtr[3], expected[3]); } } if( ! (err1 <= maxErr) || ! (err2 <= maxErr) || diff --git a/test_conformance/images/kernel_read_write/test_read_1D.cpp b/test_conformance/images/kernel_read_write/test_read_1D.cpp index e9f97340..e2e36a6f 100644 --- a/test_conformance/images/kernel_read_write/test_read_1D.cpp +++ b/test_conformance/images/kernel_read_write/test_read_1D.cpp @@ -66,8 +66,6 @@ const char *float1DKernelSource = static const char *samplerKernelArg = " sampler_t imageSampler,"; -#define ABS_ERROR( result, expected ) ( fabsf( (float)expected - (float)result ) ) - extern void read_image_pixel_float( void *imageData, image_descriptor *imageInfo, int x, int y, int z, float *outData ); template int determine_validation_error_1D( void *imagePtr, image_descriptor *imageInfo, image_sampler_data *imageSampler, @@ -522,10 +520,13 @@ int test_read_image_1D( cl_context context, cl_command_queue queue, cl_kernel ke xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f, imageSampler, expected, 0, &containsDenormals, lod ); - float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + float err2 = ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + float err3 = ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + float err4 = ABS_ERROR(resultPtr[3], expected[3]); float maxErr = 0.5; @@ -544,10 +545,13 @@ int test_read_image_1D( cl_context context, cl_command_queue queue, cl_kernel ke xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f, imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + err2 = ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + err3 = ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + err4 = ABS_ERROR(resultPtr[3], expected[3]); } } @@ -576,10 +580,14 @@ int test_read_image_1D( cl_context context, cl_command_queue queue, cl_kernel ke xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f, imageSampler, expected, 0, &containsDenormals, lod ); - float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + float err2 = ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + float err3 = ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + float err4 = + ABS_ERROR(resultPtr[3], expected[3]); float maxErr = 0.6; @@ -597,10 +605,14 @@ int test_read_image_1D( cl_context context, cl_command_queue queue, cl_kernel ke xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f, imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + err2 = ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + err3 = ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + err4 = ABS_ERROR(resultPtr[3], + expected[3]); } } if( ! (err1 <= maxErr) || ! (err2 <= maxErr) || @@ -667,10 +679,10 @@ int test_read_image_1D( cl_context context, cl_command_queue queue, cl_kernel ke xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f, imageSampler, expected, 0, &containsDenormals, lod ); - float err1 = fabsf( resultPtr[0] - expected[0] ); - float err2 = fabsf( resultPtr[1] - expected[1] ); - float err3 = fabsf( resultPtr[2] - expected[2] ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = ABS_ERROR(resultPtr[0], expected[0]); + float err2 = ABS_ERROR(resultPtr[1], expected[1]); + float err3 = ABS_ERROR(resultPtr[2], expected[2]); + float err4 = ABS_ERROR(resultPtr[3], expected[3]); // Clamp to the minimum absolute error for the format if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 0.0f; } @@ -699,10 +711,10 @@ int test_read_image_1D( cl_context context, cl_command_queue queue, cl_kernel ke xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f, imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( resultPtr[0] - expected[0] ); - err2 = fabsf( resultPtr[1] - expected[1] ); - err3 = fabsf( resultPtr[2] - expected[2] ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(resultPtr[0], expected[0]); + err2 = ABS_ERROR(resultPtr[1], expected[1]); + err3 = ABS_ERROR(resultPtr[2], expected[2]); + err4 = ABS_ERROR(resultPtr[3], expected[3]); } } @@ -731,10 +743,14 @@ int test_read_image_1D( cl_context context, cl_command_queue queue, cl_kernel ke xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f, imageSampler, expected, 0, &containsDenormals, lod ); - float err1 = fabsf( resultPtr[0] - expected[0] ); - float err2 = fabsf( resultPtr[1] - expected[1] ); - float err3 = fabsf( resultPtr[2] - expected[2] ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = + ABS_ERROR(resultPtr[0], expected[0]); + float err2 = + ABS_ERROR(resultPtr[1], expected[1]); + float err3 = + ABS_ERROR(resultPtr[2], expected[2]); + float err4 = + ABS_ERROR(resultPtr[3], expected[3]); float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN ); float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN ); float maxErr3 = MAX( maxErr * maxPixel.p[2], FLT_MIN ); @@ -756,10 +772,14 @@ int test_read_image_1D( cl_context context, cl_command_queue queue, cl_kernel ke xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f, imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( resultPtr[0] - expected[0] ); - err2 = fabsf( resultPtr[1] - expected[1] ); - err3 = fabsf( resultPtr[2] - expected[2] ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(resultPtr[0], + expected[0]); + err2 = ABS_ERROR(resultPtr[1], + expected[1]); + err3 = ABS_ERROR(resultPtr[2], + expected[2]); + err4 = ABS_ERROR(resultPtr[3], + expected[3]); } } if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) || diff --git a/test_conformance/images/kernel_read_write/test_read_1D_array.cpp b/test_conformance/images/kernel_read_write/test_read_1D_array.cpp index 28971382..eede817c 100644 --- a/test_conformance/images/kernel_read_write/test_read_1D_array.cpp +++ b/test_conformance/images/kernel_read_write/test_read_1D_array.cpp @@ -73,8 +73,6 @@ const char *floatKernelSource1DArray = static const char *samplerKernelArg = " sampler_t imageSampler,"; -#define ABS_ERROR( result, expected ) ( fabsf( (float)expected - (float)result ) ) - extern void read_image_pixel_float( void *imageData, image_descriptor *imageInfo, int x, int y, int z, float *outData ); @@ -618,10 +616,13 @@ int test_read_image_1D_array( cl_context context, cl_command_queue queue, cl_ker xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, &containsDenormals, lod ); - float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + float err2 = ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + float err3 = ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + float err4 = ABS_ERROR(resultPtr[3], expected[3]); float maxErr = 0.5; // Check if the result matches. @@ -639,10 +640,13 @@ int test_read_image_1D_array( cl_context context, cl_command_queue queue, cl_ker xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + err2 = ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + err3 = ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + err4 = ABS_ERROR(resultPtr[3], expected[3]); } } @@ -674,10 +678,14 @@ int test_read_image_1D_array( cl_context context, cl_command_queue queue, cl_ker xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, &containsDenormals, lod ); - float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + float err2 = ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + float err3 = ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + float err4 = + ABS_ERROR(resultPtr[3], expected[3]); float maxErr = 0.6; @@ -695,10 +703,14 @@ int test_read_image_1D_array( cl_context context, cl_command_queue queue, cl_ker xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + err2 = ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + err3 = ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + err4 = ABS_ERROR(resultPtr[3], + expected[3]); } } if( ! (err1 <= maxErr) || ! (err2 <= maxErr) || @@ -772,10 +784,10 @@ int test_read_image_1D_array( cl_context context, cl_command_queue queue, cl_ker xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, &containsDenormals, lod ); - float err1 = fabsf( resultPtr[0] - expected[0] ); - float err2 = fabsf( resultPtr[1] - expected[1] ); - float err3 = fabsf( resultPtr[2] - expected[2] ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = ABS_ERROR(resultPtr[0], expected[0]); + float err2 = ABS_ERROR(resultPtr[1], expected[1]); + float err3 = ABS_ERROR(resultPtr[2], expected[2]); + float err4 = ABS_ERROR(resultPtr[3], expected[3]); // Clamp to the minimum absolute error for the format if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 0.0f; } @@ -804,10 +816,10 @@ int test_read_image_1D_array( cl_context context, cl_command_queue queue, cl_ker xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( resultPtr[0] - expected[0] ); - err2 = fabsf( resultPtr[1] - expected[1] ); - err3 = fabsf( resultPtr[2] - expected[2] ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(resultPtr[0], expected[0]); + err2 = ABS_ERROR(resultPtr[1], expected[1]); + err3 = ABS_ERROR(resultPtr[2], expected[2]); + err4 = ABS_ERROR(resultPtr[3], expected[3]); } } @@ -839,10 +851,14 @@ int test_read_image_1D_array( cl_context context, cl_command_queue queue, cl_ker xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, &containsDenormals, lod ); - float err1 = fabsf( resultPtr[0] - expected[0] ); - float err2 = fabsf( resultPtr[1] - expected[1] ); - float err3 = fabsf( resultPtr[2] - expected[2] ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = + ABS_ERROR(resultPtr[0], expected[0]); + float err2 = + ABS_ERROR(resultPtr[1], expected[1]); + float err3 = + ABS_ERROR(resultPtr[2], expected[2]); + float err4 = + ABS_ERROR(resultPtr[3], expected[3]); float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN ); float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN ); float maxErr3 = MAX( maxErr * maxPixel.p[2], FLT_MIN ); @@ -864,10 +880,14 @@ int test_read_image_1D_array( cl_context context, cl_command_queue queue, cl_ker xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f, imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( resultPtr[0] - expected[0] ); - err2 = fabsf( resultPtr[1] - expected[1] ); - err3 = fabsf( resultPtr[2] - expected[2] ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(resultPtr[0], + expected[0]); + err2 = ABS_ERROR(resultPtr[1], + expected[1]); + err3 = ABS_ERROR(resultPtr[2], + expected[2]); + err4 = ABS_ERROR(resultPtr[3], + expected[3]); } } if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) || diff --git a/test_conformance/images/kernel_read_write/test_read_2D_array.cpp b/test_conformance/images/kernel_read_write/test_read_2D_array.cpp index e9212460..79420b4c 100644 --- a/test_conformance/images/kernel_read_write/test_read_2D_array.cpp +++ b/test_conformance/images/kernel_read_write/test_read_2D_array.cpp @@ -88,8 +88,6 @@ const char *float2DArrayUnnormalizedCoordKernelSource = static const char *samplerKernelArg = " sampler_t imageSampler,"; -#define ABS_ERROR( result, expected ) ( fabsf( (float)expected - (float)result ) ) - extern void read_image_pixel_float( void *imageData, image_descriptor *imageInfo, int x, int y, int z, float *outData ); template int determine_validation_error_offset_2D_array( void *imagePtr, image_descriptor *imageInfo, image_sampler_data *imageSampler, T *resultPtr, T * expected, float error, @@ -622,7 +620,8 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, &hasDenormals, lod ); - float err1 = fabsf( resultPtr[0] - expected[0] ); + float err1 = + ABS_ERROR(resultPtr[0], expected[0]); // Clamp to the minimum absolute error for the format if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN ); @@ -641,7 +640,8 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( resultPtr[0] - expected[0] ); + err1 = ABS_ERROR(resultPtr[0], + expected[0]); } } @@ -666,7 +666,8 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, &hasDenormals, lod ); - float err1 = fabsf( resultPtr[0] - expected[0] ); + float err1 = ABS_ERROR(resultPtr[0], + expected[0]); float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN ); @@ -681,7 +682,8 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker xOffsetValues[ j ], yOffsetValues[ j ], zOffsetValues[ j ], imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( resultPtr[0] - expected[0] ); + err1 = ABS_ERROR(resultPtr[0], + expected[0]); } } @@ -757,10 +759,17 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, &hasDenormals, lod ); - float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = + ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + float err2 = + ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + float err3 = + ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + float err4 = + ABS_ERROR(resultPtr[3], expected[3]); float maxErr = 0.5; if( ! (err1 <= maxErr) || ! (err2 <= maxErr) || ! (err3 <= maxErr) || ! (err4 <= maxErr) ) @@ -777,10 +786,17 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = + ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + err2 = + ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + err3 = + ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + err4 = ABS_ERROR(resultPtr[3], + expected[3]); } } @@ -805,10 +821,17 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, &hasDenormals, lod ); - float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = + ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + float err2 = + ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + float err3 = + ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + float err4 = ABS_ERROR(resultPtr[3], + expected[3]); float maxErr = 0.6; if( ! (err1 <= maxErr) || ! (err2 <= maxErr) || ! (err3 <= maxErr) || ! (err4 <= maxErr) ) @@ -824,10 +847,17 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker xOffsetValues[ j ], yOffsetValues[ j ], zOffsetValues[ j ], imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR( + sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + err2 = ABS_ERROR( + sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + err3 = ABS_ERROR( + sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + err4 = ABS_ERROR(resultPtr[3], + expected[3]); } } @@ -906,10 +936,14 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, &hasDenormals, lod ); - float err1 = fabsf( resultPtr[0] - expected[0] ); - float err2 = fabsf( resultPtr[1] - expected[1] ); - float err3 = fabsf( resultPtr[2] - expected[2] ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = + ABS_ERROR(resultPtr[0], expected[0]); + float err2 = + ABS_ERROR(resultPtr[1], expected[1]); + float err3 = + ABS_ERROR(resultPtr[2], expected[2]); + float err4 = + ABS_ERROR(resultPtr[3], expected[3]); // Clamp to the minimum absolute error for the format if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 0.0f; } @@ -937,10 +971,14 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( resultPtr[0] - expected[0] ); - err2 = fabsf( resultPtr[1] - expected[1] ); - err3 = fabsf( resultPtr[2] - expected[2] ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(resultPtr[0], + expected[0]); + err2 = ABS_ERROR(resultPtr[1], + expected[1]); + err3 = ABS_ERROR(resultPtr[2], + expected[2]); + err4 = ABS_ERROR(resultPtr[3], + expected[3]); } } @@ -965,10 +1003,14 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, &hasDenormals, lod ); - float err1 = fabsf( resultPtr[0] - expected[0] ); - float err2 = fabsf( resultPtr[1] - expected[1] ); - float err3 = fabsf( resultPtr[2] - expected[2] ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = ABS_ERROR(resultPtr[0], + expected[0]); + float err2 = ABS_ERROR(resultPtr[1], + expected[1]); + float err3 = ABS_ERROR(resultPtr[2], + expected[2]); + float err4 = ABS_ERROR(resultPtr[3], + expected[3]); float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN ); float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN ); float maxErr3 = MAX( maxErr * maxPixel.p[2], FLT_MIN ); @@ -989,10 +1031,14 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker xOffsetValues[ j ], yOffsetValues[ j ], zOffsetValues[ j ], imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( resultPtr[0] - expected[0] ); - err2 = fabsf( resultPtr[1] - expected[1] ); - err3 = fabsf( resultPtr[2] - expected[2] ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(resultPtr[0], + expected[0]); + err2 = ABS_ERROR(resultPtr[1], + expected[1]); + err3 = ABS_ERROR(resultPtr[2], + expected[2]); + err4 = ABS_ERROR(resultPtr[3], + expected[3]); } } diff --git a/test_conformance/images/kernel_read_write/test_read_3D.cpp b/test_conformance/images/kernel_read_write/test_read_3D.cpp index fe375be9..0b9e8dee 100644 --- a/test_conformance/images/kernel_read_write/test_read_3D.cpp +++ b/test_conformance/images/kernel_read_write/test_read_3D.cpp @@ -88,8 +88,6 @@ const char *float3DUnnormalizedCoordKernelSource = static const char *samplerKernelArg = " sampler_t imageSampler,"; -#define ABS_ERROR( result, expected ) ( fabsf( (float)expected - (float)result ) ) - extern void read_image_pixel_float( void *imageData, image_descriptor *imageInfo, int x, int y, int z, float *outData ); template int determine_validation_error_offset( void *imagePtr, image_descriptor *imageInfo, image_sampler_data *imageSampler, T *resultPtr, T * expected, float error, @@ -638,10 +636,17 @@ int test_read_image_3D( cl_context context, cl_command_queue queue, cl_kernel ke norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, &hasDenormals, lod ); - float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = + ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + float err2 = + ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + float err3 = + ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + float err4 = + ABS_ERROR(resultPtr[3], expected[3]); // Clamp to the minimum absolute error for the format if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 0.0f; } @@ -663,10 +668,17 @@ int test_read_image_3D( cl_context context, cl_command_queue queue, cl_kernel ke norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = + ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + err2 = + ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + err3 = + ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + err4 = ABS_ERROR(resultPtr[3], + expected[3]); } } @@ -691,10 +703,17 @@ int test_read_image_3D( cl_context context, cl_command_queue queue, cl_kernel ke norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, &hasDenormals, lod ); - float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = + ABS_ERROR(sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + float err2 = + ABS_ERROR(sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + float err3 = + ABS_ERROR(sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + float err4 = ABS_ERROR(resultPtr[3], + expected[3]); float maxErr = 0.6; if( ! (err1 <= maxErr) || ! (err2 <= maxErr) || ! (err3 <= maxErr) || ! (err4 <= maxErr) ) @@ -710,10 +729,17 @@ int test_read_image_3D( cl_context context, cl_command_queue queue, cl_kernel ke xOffsetValues[ j ], yOffsetValues[ j ], zOffsetValues[ j ], imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); - err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); - err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR( + sRGBmap(resultPtr[0]), + sRGBmap(expected[0])); + err2 = ABS_ERROR( + sRGBmap(resultPtr[1]), + sRGBmap(expected[1])); + err3 = ABS_ERROR( + sRGBmap(resultPtr[2]), + sRGBmap(expected[2])); + err4 = ABS_ERROR(resultPtr[3], + expected[3]); } } @@ -792,10 +818,14 @@ int test_read_image_3D( cl_context context, cl_command_queue queue, cl_kernel ke norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, &hasDenormals, lod ); - float err1 = fabsf( resultPtr[0] - expected[0] ); - float err2 = fabsf( resultPtr[1] - expected[1] ); - float err3 = fabsf( resultPtr[2] - expected[2] ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = + ABS_ERROR(resultPtr[0], expected[0]); + float err2 = + ABS_ERROR(resultPtr[1], expected[1]); + float err3 = + ABS_ERROR(resultPtr[2], expected[2]); + float err4 = + ABS_ERROR(resultPtr[3], expected[3]); // Clamp to the minimum absolute error for the format if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 0.0f; } @@ -823,10 +853,14 @@ int test_read_image_3D( cl_context context, cl_command_queue queue, cl_kernel ke norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( resultPtr[0] - expected[0] ); - err2 = fabsf( resultPtr[1] - expected[1] ); - err3 = fabsf( resultPtr[2] - expected[2] ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(resultPtr[0], + expected[0]); + err2 = ABS_ERROR(resultPtr[1], + expected[1]); + err3 = ABS_ERROR(resultPtr[2], + expected[2]); + err4 = ABS_ERROR(resultPtr[3], + expected[3]); } } @@ -851,10 +885,14 @@ int test_read_image_3D( cl_context context, cl_command_queue queue, cl_kernel ke norm_offset_x, norm_offset_y, norm_offset_z, imageSampler, expected, 0, &hasDenormals, lod ); - float err1 = fabsf( resultPtr[0] - expected[0] ); - float err2 = fabsf( resultPtr[1] - expected[1] ); - float err3 = fabsf( resultPtr[2] - expected[2] ); - float err4 = fabsf( resultPtr[3] - expected[3] ); + float err1 = ABS_ERROR(resultPtr[0], + expected[0]); + float err2 = ABS_ERROR(resultPtr[1], + expected[1]); + float err3 = ABS_ERROR(resultPtr[2], + expected[2]); + float err4 = ABS_ERROR(resultPtr[3], + expected[3]); float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN ); float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN ); float maxErr3 = MAX( maxErr * maxPixel.p[2], FLT_MIN ); @@ -875,10 +913,14 @@ int test_read_image_3D( cl_context context, cl_command_queue queue, cl_kernel ke xOffsetValues[ j ], yOffsetValues[ j ], zOffsetValues[ j ], imageSampler, expected, 0, NULL, lod ); - err1 = fabsf( resultPtr[0] - expected[0] ); - err2 = fabsf( resultPtr[1] - expected[1] ); - err3 = fabsf( resultPtr[2] - expected[2] ); - err4 = fabsf( resultPtr[3] - expected[3] ); + err1 = ABS_ERROR(resultPtr[0], + expected[0]); + err2 = ABS_ERROR(resultPtr[1], + expected[1]); + err3 = ABS_ERROR(resultPtr[2], + expected[2]); + err4 = ABS_ERROR(resultPtr[3], + expected[3]); } } @@ -1311,7 +1353,3 @@ int test_read_image_set_3D( cl_device_id device, cl_context context, cl_command_ return 0; } - - - - diff --git a/test_conformance/math_brute_force/main.cpp b/test_conformance/math_brute_force/main.cpp index 878c2f34..1e33b95a 100644 --- a/test_conformance/math_brute_force/main.cpp +++ b/test_conformance/math_brute_force/main.cpp @@ -15,10 +15,9 @@ // #include "Utility.h" -#include -#include - -#include +#include +#include +#include #include #include "FunctionList.h" #include "Sleep.h" @@ -106,24 +105,6 @@ cl_device_fp_config gDoubleCapabilities = 0; int gWimpyReductionFactor = 32; int gWimpyBufferSize = BUFFER_SIZE; int gVerboseBruteForce = 0; -#if defined( __APPLE__ ) -int gHasBasicDouble = 0; -char* gBasicDoubleFuncs[] = { - "add", - "assignment", - "divide", - "isequal", - "isgreater", - "isgreaterequal", - "isless", - "islessequal", - "isnotequal", - "multiply", - "sqrt", - "subtract" }; -size_t gNumBasicDoubleFuncs = sizeof(gBasicDoubleFuncs)/sizeof(char*); -#endif - static int ParseArgs( int argc, const char **argv ); static void PrintUsage( void ); @@ -255,44 +236,6 @@ int doTest( const char* name ) //Re-enable testing fast-relaxed-math mode gTestFastRelaxed = testFastRelaxedTmp; } - -#if defined( __APPLE__ ) - { - if( gHasBasicDouble && NULL != func_data->vtbl_ptr->DoubleTestFunc && NULL != func_data->dfunc.p) - { - //Disable fast-relaxed-math for double precision floating-point - int testFastRelaxedTmp = gTestFastRelaxed; - gTestFastRelaxed = 0; - - int isBasicTest = 0; - for( size_t j = 0; j < gNumBasicDoubleFuncs; j++ ) { - if( 0 == strcmp(gBasicDoubleFuncs[j], func_data->name ) ) { - isBasicTest = 1; - break; - } - } - if (isBasicTest) { - gTestCount++; - if( gTestFloat ) - vlog( " " ); - if( func_data->vtbl_ptr->DoubleTestFunc( func_data, gMTdata ) ) - { - gFailCount++; - error++; - if( gStopOnError ) - { - gTestFastRelaxed = testFastRelaxedTmp; - gSkipRestOfTests = true; - return error; - } - } - } - - //Re-enable testing fast-relaxed-math mode - gTestFastRelaxed = testFastRelaxedTmp; - } - } -#endif } return error; @@ -1179,20 +1122,20 @@ test_status InitCL( cl_device_id device ) if( DOUBLE_REQUIRED_FEATURES != (gDoubleCapabilities & DOUBLE_REQUIRED_FEATURES) ) { - char list[300] = ""; - if( 0 == (gDoubleCapabilities & CL_FP_FMA) ) - strncat( list, "CL_FP_FMA, ", sizeof( list )-1 ); + std::string list; + if (0 == (gDoubleCapabilities & CL_FP_FMA)) list += "CL_FP_FMA, "; if( 0 == (gDoubleCapabilities & CL_FP_ROUND_TO_NEAREST) ) - strncat( list, "CL_FP_ROUND_TO_NEAREST, ", sizeof( list )-1 ); + list += "CL_FP_ROUND_TO_NEAREST, "; if( 0 == (gDoubleCapabilities & CL_FP_ROUND_TO_ZERO) ) - strncat( list, "CL_FP_ROUND_TO_ZERO, ", sizeof( list )-1 ); + list += "CL_FP_ROUND_TO_ZERO, "; if( 0 == (gDoubleCapabilities & CL_FP_ROUND_TO_INF) ) - strncat( list, "CL_FP_ROUND_TO_INF, ", sizeof( list )-1 ); + list += "CL_FP_ROUND_TO_INF, "; if( 0 == (gDoubleCapabilities & CL_FP_INF_NAN) ) - strncat( list, "CL_FP_INF_NAN, ", sizeof( list )-1 ); + list += "CL_FP_INF_NAN, "; if( 0 == (gDoubleCapabilities & CL_FP_DENORM) ) - strncat( list, "CL_FP_DENORM, ", sizeof( list )-1 ); - vlog_error( "ERROR: required double features are missing: %s\n", list ); + list += "CL_FP_DENORM, "; + vlog_error("ERROR: required double features are missing: %s\n", + list.c_str()); return TEST_FAIL; } @@ -1201,43 +1144,6 @@ test_status InitCL( cl_device_id device ) return TEST_FAIL; #endif } -#if defined( __APPLE__ ) - else if(is_extension_available(gDevice, "cl_APPLE_fp64_basic_ops" )) - { - gHasBasicDouble ^= 1; - -#if defined( CL_DEVICE_DOUBLE_FP_CONFIG ) - if( (error = clGetDeviceInfo(gDevice, CL_DEVICE_DOUBLE_FP_CONFIG, sizeof(gDoubleCapabilities), &gDoubleCapabilities, NULL))) - { - vlog_error( "ERROR: Unable to get device CL_DEVICE_DOUBLE_FP_CONFIG. (%d)\n", error ); - return TEST_FAIL; - } - - if( DOUBLE_REQUIRED_FEATURES != (gDoubleCapabilities & DOUBLE_REQUIRED_FEATURES) ) - { - char list[300] = ""; - if( 0 == (gDoubleCapabilities & CL_FP_FMA) ) - strncat( list, "CL_FP_FMA, ", sizeof( list ) ); - if( 0 == (gDoubleCapabilities & CL_FP_ROUND_TO_NEAREST) ) - strncat( list, "CL_FP_ROUND_TO_NEAREST, ", sizeof( list ) ); - if( 0 == (gDoubleCapabilities & CL_FP_ROUND_TO_ZERO) ) - strncat( list, "CL_FP_ROUND_TO_ZERO, ", sizeof( list ) ); - if( 0 == (gDoubleCapabilities & CL_FP_ROUND_TO_INF) ) - strncat( list, "CL_FP_ROUND_TO_INF, ", sizeof( list ) ); - if( 0 == (gDoubleCapabilities & CL_FP_INF_NAN) ) - strncat( list, "CL_FP_INF_NAN, ", sizeof( list ) ); - if( 0 == (gDoubleCapabilities & CL_FP_DENORM) ) - strncat( list, "CL_FP_DENORM, ", sizeof( list ) ); - vlog_error( "ERROR: required double features are missing: %s\n", list ); - - return TEST_FAIL; - } -#else - vlog_error( "FAIL: device says it supports cl_khr_fp64 but CL_DEVICE_DOUBLE_FP_CONFIG is not in the headers!\n" ); - return TEST_FAIL; -#endif - } -#endif /* __APPLE__ */ configSize = sizeof( gDeviceFrequency ); if( (error = clGetDeviceInfo( gDevice, CL_DEVICE_MAX_CLOCK_FREQUENCY, configSize, &gDeviceFrequency, NULL )) ) @@ -1411,9 +1317,6 @@ test_status InitCL( cl_device_id device ) vlog( "\t\t can not accurately represent the right result to an accuracy closer\n" ); vlog( "\t\t than half an ulp. See comments in Bruteforce_Ulp_Error_Double() for more details.\n\n" ); } -#if defined( __APPLE__ ) - vlog( "\tTesting basic double precision? %s\n", no_yes[0 != gHasBasicDouble] ); -#endif vlog( "\tIs Embedded? %s\n", no_yes[0 != gIsEmbedded] ); if( gIsEmbedded ) diff --git a/test_conformance/math_brute_force/reference_math.cpp b/test_conformance/math_brute_force/reference_math.cpp index b6df2436..99c8eb37 100644 --- a/test_conformance/math_brute_force/reference_math.cpp +++ b/test_conformance/math_brute_force/reference_math.cpp @@ -37,39 +37,6 @@ #define EVALUATE( x ) x #define CONCATENATE(x, y) x ## EVALUATE(y) - -// Declare Classification macros for non-C99 platforms -#ifndef isinf - #define isinf(x) ( sizeof (x) == sizeof(float ) ? fabsf(x) == INFINITY \ - : sizeof (x) == sizeof(double) ? fabs(x) == INFINITY \ - : fabsl(x) == INFINITY) -#endif - -#ifndef isfinite - #define isfinite(x) ( sizeof (x) == sizeof(float ) ? fabsf(x) < INFINITY \ - : sizeof (x) == sizeof(double) ? fabs(x) < INFINITY \ - : fabsl(x) < INFINITY) -#endif - -#ifndef isnan - #define isnan(_a) ( (_a) != (_a) ) -#endif - -#ifdef __MINGW32__ - #undef isnormal -#endif - -#ifndef isnormal - #define isnormal(x) ( sizeof (x) == sizeof(float ) ? (fabsf(x) < INFINITY && fabsf(x) >= FLT_MIN) \ - : sizeof (x) == sizeof(double) ? (fabs(x) < INFINITY && fabs(x) >= DBL_MIN) \ - : (fabsl(x) < INFINITY && fabsl(x) >= LDBL_MIN) ) -#endif - -#ifndef islessgreater - // Note: Non-C99 conformant. This will trigger floating point exceptions. We don't care about that here. - #define islessgreater( _x, _y ) ( (_x) < (_y) || (_x) > (_y) ) -#endif - #pragma STDC FP_CONTRACT OFF static void __log2_ep(double *hi, double *lo, double x); diff --git a/test_conformance/spir/datagen.h b/test_conformance/spir/datagen.h index 739fb5c6..cf620466 100644 --- a/test_conformance/spir/datagen.h +++ b/test_conformance/spir/datagen.h @@ -274,20 +274,16 @@ public: if (CL_SUCCESS != error) throw Exceptions::TestError("clGetSupportedImageFormats failed\n", error); - std::auto_ptr supportedFormats(new cl_image_format[actualNumFormats]); - error = clGetSupportedImageFormats( - context, - flags, - m_desc.image_type, - actualNumFormats, - supportedFormats.get(), - NULL); + std::vector supportedFormats(actualNumFormats); + error = clGetSupportedImageFormats(context, flags, m_desc.image_type, + actualNumFormats, + supportedFormats.data(), NULL); if (CL_SUCCESS != error) throw Exceptions::TestError("clGetSupportedImageFormats failed\n", error); for (size_t i=0; i& versions) cl_int err; size_t size = 0; - if (err = clGetDeviceInfo(device, CL_DEVICE_SPIR_VERSIONS, sizeof(version), - (void*)version, &size)) + if ((err = clGetDeviceInfo(device, CL_DEVICE_SPIR_VERSIONS, sizeof(version), + (void *)version, &size))) { log_error( "Error: failed to obtain SPIR version at %s:%d (err = %d)\n", __FILE__, __LINE__, err ); diff --git a/test_conformance/spir/run_build_test.cpp b/test_conformance/spir/run_build_test.cpp index b13282da..cec2d275 100644 --- a/test_conformance/spir/run_build_test.cpp +++ b/test_conformance/spir/run_build_test.cpp @@ -268,7 +268,7 @@ static bool run_test(cl_context context, cl_command_queue queue, cl_program clpr { WorkSizeInfo ws; TestResult cl_result; - std::auto_ptr bc_result; + std::unique_ptr bc_result; // first, run the single CL test { // make sure that the kernel will be released before the program