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

Enable -Werror for GCC/Clang builds (#786)

Browse Source 
* 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
This commit is contained in:
James Price
2020-05-27 14:13:11 -04:00
committed by GitHub
parent 094cc04e16
commit 944b0a8178
27 changed files with 398 additions and 367 deletions
Download Patch File Download Diff File Expand all files Collapse all files

20
CMakeLists.txt
View File

@@ -124,6 +124,7 @@ endif (GL_IS_SUPPORTED AND CLConform_GL_LIBRARIES_DIR)
include(CheckFunctionExists) include(CheckFunctionExists)
include(CheckIncludeFiles) include(CheckIncludeFiles)
include(CheckCXXCompilerFlag)
if(CMAKE_SYSTEM_PROCESSOR MATCHES "^(arm.*|ARM.*)") if(CMAKE_SYSTEM_PROCESSOR MATCHES "^(arm.*|ARM.*)")
set(CLConform_TARGET_ARCH ARM) set(CLConform_TARGET_ARCH ARM)
@@ -139,9 +140,24 @@ if(NOT DEFINED CLConform_TARGET_ARCH)
message (FATAL_ERROR "Target architecture not recognised. Exiting.") message (FATAL_ERROR "Target architecture not recognised. Exiting.")
endif() 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") if(CMAKE_COMPILER_IS_GNUCC OR "${CMAKE_CXX_COMPILER_ID}" MATCHES "(Apple)?Clang")
set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -Wno-narrowing") add_cxx_flag_if_supported(-Wno-narrowing)
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -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, # -msse -mfpmath=sse to force gcc to use sse for float math,
# avoiding excess precision problems that cause tests like int2float # avoiding excess precision problems that cause tests like int2float
# to falsely fail. -ffloat-store also works, but WG suggested # to falsely fail. -ffloat-store also works, but WG suggested

3
test_common/harness/imageHelpers.cpp
View File

@@ -899,8 +899,7 @@ int get_format_min_int( cl_image_format *format )
case CL_UNORM_INT_101010: case CL_UNORM_INT_101010:
return 0; return 0;
case CL_HALF_FLOAT: case CL_HALF_FLOAT: return -(1 << 10);
return -1<<10;
#ifdef CL_SFIXED14_APPLE #ifdef CL_SFIXED14_APPLE
case CL_SFIXED14_APPLE: case CL_SFIXED14_APPLE:

2
test_conformance/api/test_null_buffer_arg.cpp
View File

@@ -62,7 +62,7 @@ static int test_setargs_and_execution(cl_command_queue queue, cl_kernel kernel,
unsigned int i; unsigned int i;
cl_int status; cl_int status;
char *typestr; const char *typestr;
if (type == NON_NULL_PATH) { if (type == NON_NULL_PATH) {
status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &test_buf); status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &test_buf);

6
test_conformance/basic/test_work_item_functions.cpp
View File

@@ -130,7 +130,7 @@ int test_work_item_functions(cl_device_id deviceID, cl_context context, cl_comma
free_mtdata(d); free_mtdata(d);
return -1; 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", 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 ] ); (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); free_mtdata(d);
return -1; 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", 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 ] ); (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); free_mtdata(d);
return -1; 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", 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 ] ); (int)j, (int)dim, (int)groupCount, (int)testData[ q ].groupID[ j ] );

12
test_conformance/buffers/test_buffer_migrate.cpp
View File

@@ -69,6 +69,7 @@ static cl_int migrateMemObject(enum migrations migrate, cl_command_queue *queues
// Choose a random set of flags // 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));; flags[i] = (cl_mem_migration_flags)(genrand_int32(d) & (CL_MIGRATE_MEM_OBJECT_HOST | CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED));;
break; 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) { 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."); 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. // Build the kernel program.
if (err = create_single_kernel_helper(ctx, &program, &kernel, 1, &buffer_migrate_kernel_code, "test_buffer_migrate")) { if ((err = create_single_kernel_helper(ctx, &program, &kernel, 1,
print_error(err, "Failed creating kernel."); &buffer_migrate_kernel_code,
failed = 1; "test_buffer_migrate")))
goto cleanup; {
print_error(err, "Failed creating kernel.");
failed = 1;
goto cleanup;
} }
num_devices_limited = num_devices; num_devices_limited = num_devices;

12
test_conformance/buffers/test_image_migrate.cpp
View File

@@ -80,6 +80,7 @@ static cl_int migrateMemObject(enum migrations migrate, cl_command_queue *queues
// Choose a random set of flags // 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)); flags[i] = (cl_mem_migration_flags)(genrand_int32(d) & (CL_MIGRATE_MEM_OBJECT_HOST | CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED));
break; break;
default: log_error("Unhandled migration type: %d\n", migrate); return -1;
} }
if ((err = clEnqueueMigrateMemObjects(queues[j], 1, (const cl_mem *)(&mem_objects[i]), if ((err = clEnqueueMigrateMemObjects(queues[j], 1, (const cl_mem *)(&mem_objects[i]),
flags[i], 0, NULL, NULL)) != CL_SUCCESS) { 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. // Build the kernel program.
if (err = create_single_kernel_helper(ctx, &program, &kernel, 1, &image_migrate_kernel_code, "test_image_migrate")) { if ((err = create_single_kernel_helper(ctx, &program, &kernel, 1,
print_error(err, "Failed creating kernel."); &image_migrate_kernel_code,
failed = 1; "test_image_migrate")))
goto cleanup; {
print_error(err, "Failed creating kernel.");
failed = 1;
goto cleanup;
} }
// Create sampler. // Create sampler.

2
test_conformance/compiler/test_image_macro.cpp
View File

@@ -57,7 +57,7 @@ int test_image_macro(cl_device_id deviceID, cl_context context, cl_command_queue
return status; 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); status = create_single_kernel_helper_create_program(context, &program, 1, (const char**)&image_supported_source);

2
test_conformance/computeinfo/main.cpp
View File

@@ -96,7 +96,7 @@ typedef struct _extensions extensions_t;
// first is greater). // first is greater).
int vercmp(version_t a, version_t b) 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; return -1;
} }

3
test_conformance/half/cl_utils.h
View File

@@ -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); 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 ); 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 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; } static inline double as_double(cl_ulong u) { union { cl_ulong u; double d; }v; v.u = u; return v.d; }

8
test_conformance/images/clCopyImage/test_copy_2D_2D_array.cpp
View File

@@ -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[ 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[ 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; 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 ) 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; 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; 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[ 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[ 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; 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 ) 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; 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; 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! // Go for it!
retCode = test_copy_image_generic( context, queue, srcImageInfo, dstImageInfo, sourcePos, destPos, regionSize, d ); retCode = test_copy_image_generic( context, queue, srcImageInfo, dstImageInfo, sourcePos, destPos, regionSize, d );

8
test_conformance/images/clCopyImage/test_copy_2D_3D.cpp
View File

@@ -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[ 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; 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 ) 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; 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; 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[ 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[ 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; 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 ) 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; 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; 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! // Go for it!
retCode = test_copy_image_generic( context, queue, srcImageInfo, dstImageInfo, sourcePos, destPos, regionSize, d ); retCode = test_copy_image_generic( context, queue, srcImageInfo, dstImageInfo, sourcePos, destPos, regionSize, d );

2
test_conformance/images/kernel_image_methods/test_1D.cpp
View File

@@ -20,7 +20,7 @@
extern bool gDebugTrace, gTestSmallImages, gTestMaxImages, gDeviceLt20; extern bool gDebugTrace, gTestSmallImages, gTestMaxImages, gDeviceLt20;
typedef struct image_kernel_data struct image_kernel_data
{ {
cl_int width; cl_int width;
cl_int channelType; cl_int channelType;

2
test_conformance/images/kernel_image_methods/test_1D_array.cpp
View File

@@ -20,7 +20,7 @@
extern bool gDebugTrace, gTestSmallImages, gTestMaxImages, gDeviceLt20; extern bool gDebugTrace, gTestSmallImages, gTestMaxImages, gDeviceLt20;
typedef struct image_kernel_data struct image_kernel_data
{ {
cl_int width; cl_int width;
cl_int arraySize; cl_int arraySize;

2
test_conformance/images/kernel_image_methods/test_2D.cpp
View File

@@ -20,7 +20,7 @@
extern bool gDebugTrace, gTestSmallImages, gTestMaxImages, gDeviceLt20; extern bool gDebugTrace, gTestSmallImages, gTestMaxImages, gDeviceLt20;
typedef struct image_kernel_data struct image_kernel_data
{ {
cl_int width; cl_int width;
cl_int height; cl_int height;

2
test_conformance/images/kernel_image_methods/test_2D_array.cpp
View File

@@ -20,7 +20,7 @@
extern bool gDebugTrace, gTestSmallImages, gTestMaxImages, gDeviceLt20; extern bool gDebugTrace, gTestSmallImages, gTestMaxImages, gDeviceLt20;
typedef struct image_kernel_data struct image_kernel_data
{ {
cl_int width; cl_int width;
cl_int height; cl_int height;

5
test_conformance/images/kernel_image_methods/test_loops.cpp
View File

@@ -111,7 +111,8 @@ int test_image_set( cl_device_id device, cl_context context, cl_command_queue qu
gDeviceLt20 = true; gDeviceLt20 = true;
} }
if (version_check = (version < Version(1,2))) { if ((version_check = (version < Version(1, 2))))
{
switch (imageType) { switch (imageType) {
case CL_MEM_OBJECT_IMAGE1D: case CL_MEM_OBJECT_IMAGE1D:
test_missing_feature(version_check, "image_1D"); 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: case CL_MEM_OBJECT_IMAGE2D_ARRAY:
test_missing_feature(version_check, "image_2D_array"); test_missing_feature(version_check, "image_2D_array");
} }
} }
int ret = 0; int ret = 0;
ret += test_image_type( device, context, queue, imageType, CL_MEM_READ_ONLY ); ret += test_image_type( device, context, queue, imageType, CL_MEM_READ_ONLY );

2
test_conformance/images/kernel_read_write/test_common.h
View File

@@ -1,5 +1,7 @@
#include "../testBase.h" #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); extern cl_sampler create_sampler(cl_context context, image_sampler_data *sdata, bool test_mipmaps, cl_int *error);

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

@@ -85,8 +85,6 @@ static const char *lodOffsetSource =
static const char *offsetSource = static const char *offsetSource =
" int offset = tidY*get_image_width(input) + tidX;\n"; " 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, extern void read_image_pixel_float( void *imageData, image_descriptor *imageInfo,
int x, int y, int z, float *outData ); int x, int y, int z, float *outData );
template <class T> int determine_validation_error( void *imagePtr, image_descriptor *imageInfo, image_sampler_data *imageSampler, template <class T> 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.0f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, &containsDenormals ); 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 // Clamp to the minimum absolute error for the format
if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; }
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN ); 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, NULL ); 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, &containsDenormals ); 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 ); 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, NULL ); imageSampler, expected, 0, NULL );
err1 = fabsf( resultPtr[0] - expected[0] ); err1 = ABS_ERROR(resultPtr[0], expected[0]);
} }
} }
if( ! (err1 <= maxErr1) ) 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.0f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, &containsDenormals ); imageSampler, expected, 0, &containsDenormals );
float err1 = fabsf( resultPtr[0] - expected[0] ); float err1 = ABS_ERROR(resultPtr[0], expected[0]);
float err2 = fabsf( resultPtr[1] - expected[1] ); float err2 = ABS_ERROR(resultPtr[1], expected[1]);
float err3 = fabsf( resultPtr[2] - expected[2] ); float err3 = ABS_ERROR(resultPtr[2], expected[2]);
float err4 = fabsf( resultPtr[3] - expected[3] ); float err4 = ABS_ERROR(resultPtr[3], expected[3]);
// Clamp to the minimum absolute error for the format // Clamp to the minimum absolute error for the format
if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; }
if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, NULL ); imageSampler, expected, 0, NULL );
err1 = fabsf( resultPtr[0] - expected[0] ); err1 = ABS_ERROR(resultPtr[0], expected[0]);
err2 = fabsf( resultPtr[1] - expected[1] ); err2 = ABS_ERROR(resultPtr[1], expected[1]);
err3 = fabsf( resultPtr[2] - expected[2] ); err3 = ABS_ERROR(resultPtr[2], expected[2]);
err4 = fabsf( resultPtr[3] - expected[3] ); 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, &containsDenormals ); imageSampler, expected, 0, &containsDenormals );
float err1 = fabsf( resultPtr[0] - expected[0] ); float err1 = ABS_ERROR(resultPtr[0], expected[0]);
float err2 = fabsf( resultPtr[1] - expected[1] ); float err2 = ABS_ERROR(resultPtr[1], expected[1]);
float err3 = fabsf( resultPtr[2] - expected[2] ); float err3 = ABS_ERROR(resultPtr[2], expected[2]);
float err4 = fabsf( resultPtr[3] - expected[3] ); float err4 = ABS_ERROR(resultPtr[3], expected[3]);
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN ); float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN ); float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN );
float maxErr3 = MAX( maxErr * maxPixel.p[2], 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, NULL ); imageSampler, expected, 0, NULL );
err1 = fabsf( resultPtr[0] - expected[0] ); err1 = ABS_ERROR(resultPtr[0], expected[0]);
err2 = fabsf( resultPtr[1] - expected[1] ); err2 = ABS_ERROR(resultPtr[1], expected[1]);
err3 = fabsf( resultPtr[2] - expected[2] ); err3 = ABS_ERROR(resultPtr[2], expected[2]);
err4 = fabsf( resultPtr[3] - expected[3] ); err4 = ABS_ERROR(resultPtr[3], expected[3]);
} }
} }
if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) || 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, maxPixel = sample_image_pixel_float_offset( imageValues, imageInfo,
xOffsetValues[ j ], yOffsetValues[ j ], 0.0f, norm_offset_x, norm_offset_y, 0.0f, xOffsetValues[ j ], yOffsetValues[ j ], 0.0f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, &containsDenormals ); imageSampler, expected, 0, &containsDenormals );
float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); float err1 = ABS_ERROR(sRGBmap(resultPtr[0]),
float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); sRGBmap(expected[0]));
float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); float err2 = ABS_ERROR(sRGBmap(resultPtr[1]),
float err4 = fabsf( resultPtr[3] - expected[3] ); 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; float maxErr = 0.5;
// Check if the result matches. // 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, NULL ); imageSampler, expected, 0, NULL );
err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); err1 = ABS_ERROR(sRGBmap(resultPtr[0]),
err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); sRGBmap(expected[0]));
err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); err2 = ABS_ERROR(sRGBmap(resultPtr[1]),
err4 = fabsf( resultPtr[3] - expected[3] ); 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, &containsDenormals ); imageSampler, expected, 0, &containsDenormals );
float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); float err1 = ABS_ERROR(sRGBmap(resultPtr[0]),
float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); sRGBmap(expected[0]));
float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); float err2 = ABS_ERROR(sRGBmap(resultPtr[1]),
float err4 = fabsf( resultPtr[3] - expected[3] ); 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; float maxErr = 0.6;
if( ! (err1 <= maxErr) || ! (err2 <= maxErr) || 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, NULL ); imageSampler, expected, 0, NULL );
err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); err1 = ABS_ERROR(sRGBmap(resultPtr[0]),
err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); sRGBmap(expected[0]));
err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); err2 = ABS_ERROR(sRGBmap(resultPtr[1]),
err4 = fabsf( resultPtr[3] - expected[3] ); sRGBmap(expected[1]));
err3 = ABS_ERROR(sRGBmap(resultPtr[2]),
sRGBmap(expected[2]));
err4 = ABS_ERROR(resultPtr[3], expected[3]);
} }
} }
if( ! (err1 <= maxErr) || ! (err2 <= maxErr) || if( ! (err1 <= maxErr) || ! (err2 <= maxErr) ||

88
test_conformance/images/kernel_read_write/test_read_1D.cpp
View File

@@ -66,8 +66,6 @@ const char *float1DKernelSource =
static const char *samplerKernelArg = " sampler_t imageSampler,"; 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, extern void read_image_pixel_float( void *imageData, image_descriptor *imageInfo,
int x, int y, int z, float *outData ); int x, int y, int z, float *outData );
template <class T> int determine_validation_error_1D( void *imagePtr, image_descriptor *imageInfo, image_sampler_data *imageSampler, template <class T> 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, xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f,
imageSampler, expected, 0, &containsDenormals, lod ); imageSampler, expected, 0, &containsDenormals, lod );
float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); float err1 = ABS_ERROR(sRGBmap(resultPtr[0]),
float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); sRGBmap(expected[0]));
float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); float err2 = ABS_ERROR(sRGBmap(resultPtr[1]),
float err4 = fabsf( resultPtr[3] - expected[3] ); 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; 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, xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f,
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); err1 = ABS_ERROR(sRGBmap(resultPtr[0]),
err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); sRGBmap(expected[0]));
err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); err2 = ABS_ERROR(sRGBmap(resultPtr[1]),
err4 = fabsf( resultPtr[3] - expected[3] ); 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, xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f,
imageSampler, expected, 0, &containsDenormals, lod ); imageSampler, expected, 0, &containsDenormals, lod );
float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); float err1 = ABS_ERROR(sRGBmap(resultPtr[0]),
float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); sRGBmap(expected[0]));
float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); float err2 = ABS_ERROR(sRGBmap(resultPtr[1]),
float err4 = fabsf( resultPtr[3] - expected[3] ); 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; 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, xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f,
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); err1 = ABS_ERROR(sRGBmap(resultPtr[0]),
err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); sRGBmap(expected[0]));
err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); err2 = ABS_ERROR(sRGBmap(resultPtr[1]),
err4 = fabsf( resultPtr[3] - expected[3] ); sRGBmap(expected[1]));
err3 = ABS_ERROR(sRGBmap(resultPtr[2]),
sRGBmap(expected[2]));
err4 = ABS_ERROR(resultPtr[3],
expected[3]);
} }
} }
if( ! (err1 <= maxErr) || ! (err2 <= maxErr) || 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, xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f,
imageSampler, expected, 0, &containsDenormals, lod ); imageSampler, expected, 0, &containsDenormals, lod );
float err1 = fabsf( resultPtr[0] - expected[0] ); float err1 = ABS_ERROR(resultPtr[0], expected[0]);
float err2 = fabsf( resultPtr[1] - expected[1] ); float err2 = ABS_ERROR(resultPtr[1], expected[1]);
float err3 = fabsf( resultPtr[2] - expected[2] ); float err3 = ABS_ERROR(resultPtr[2], expected[2]);
float err4 = fabsf( resultPtr[3] - expected[3] ); float err4 = ABS_ERROR(resultPtr[3], expected[3]);
// Clamp to the minimum absolute error for the format // Clamp to the minimum absolute error for the format
if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; }
if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 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, xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f,
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( resultPtr[0] - expected[0] ); err1 = ABS_ERROR(resultPtr[0], expected[0]);
err2 = fabsf( resultPtr[1] - expected[1] ); err2 = ABS_ERROR(resultPtr[1], expected[1]);
err3 = fabsf( resultPtr[2] - expected[2] ); err3 = ABS_ERROR(resultPtr[2], expected[2]);
err4 = fabsf( resultPtr[3] - expected[3] ); 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, xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f,
imageSampler, expected, 0, &containsDenormals, lod ); imageSampler, expected, 0, &containsDenormals, lod );
float err1 = fabsf( resultPtr[0] - expected[0] ); float err1 =
float err2 = fabsf( resultPtr[1] - expected[1] ); ABS_ERROR(resultPtr[0], expected[0]);
float err3 = fabsf( resultPtr[2] - expected[2] ); float err2 =
float err4 = fabsf( resultPtr[3] - expected[3] ); 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 maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN ); float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN );
float maxErr3 = MAX( maxErr * maxPixel.p[2], 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, xOffsetValues[ j ], 0.0f, 0.0f, norm_offset_x, 0.0f, 0.0f,
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( resultPtr[0] - expected[0] ); err1 = ABS_ERROR(resultPtr[0],
err2 = fabsf( resultPtr[1] - expected[1] ); expected[0]);
err3 = fabsf( resultPtr[2] - expected[2] ); err2 = ABS_ERROR(resultPtr[1],
err4 = fabsf( resultPtr[3] - expected[3] ); expected[1]);
err3 = ABS_ERROR(resultPtr[2],
expected[2]);
err4 = ABS_ERROR(resultPtr[3],
expected[3]);
} }
} }
if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) || if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) ||

88
test_conformance/images/kernel_read_write/test_read_1D_array.cpp
View File

@@ -73,8 +73,6 @@ const char *floatKernelSource1DArray =
static const char *samplerKernelArg = " sampler_t imageSampler,"; 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, extern void read_image_pixel_float( void *imageData, image_descriptor *imageInfo,
int x, int y, int z, float *outData ); 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, &containsDenormals, lod ); imageSampler, expected, 0, &containsDenormals, lod );
float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); float err1 = ABS_ERROR(sRGBmap(resultPtr[0]),
float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); sRGBmap(expected[0]));
float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); float err2 = ABS_ERROR(sRGBmap(resultPtr[1]),
float err4 = fabsf( resultPtr[3] - expected[3] ); 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; float maxErr = 0.5;
// Check if the result matches. // 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); err1 = ABS_ERROR(sRGBmap(resultPtr[0]),
err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); sRGBmap(expected[0]));
err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); err2 = ABS_ERROR(sRGBmap(resultPtr[1]),
err4 = fabsf( resultPtr[3] - expected[3] ); 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, &containsDenormals, lod ); imageSampler, expected, 0, &containsDenormals, lod );
float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); float err1 = ABS_ERROR(sRGBmap(resultPtr[0]),
float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); sRGBmap(expected[0]));
float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); float err2 = ABS_ERROR(sRGBmap(resultPtr[1]),
float err4 = fabsf( resultPtr[3] - expected[3] ); 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; 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); err1 = ABS_ERROR(sRGBmap(resultPtr[0]),
err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); sRGBmap(expected[0]));
err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); err2 = ABS_ERROR(sRGBmap(resultPtr[1]),
err4 = fabsf( resultPtr[3] - expected[3] ); sRGBmap(expected[1]));
err3 = ABS_ERROR(sRGBmap(resultPtr[2]),
sRGBmap(expected[2]));
err4 = ABS_ERROR(resultPtr[3],
expected[3]);
} }
} }
if( ! (err1 <= maxErr) || ! (err2 <= maxErr) || 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, &containsDenormals, lod ); imageSampler, expected, 0, &containsDenormals, lod );
float err1 = fabsf( resultPtr[0] - expected[0] ); float err1 = ABS_ERROR(resultPtr[0], expected[0]);
float err2 = fabsf( resultPtr[1] - expected[1] ); float err2 = ABS_ERROR(resultPtr[1], expected[1]);
float err3 = fabsf( resultPtr[2] - expected[2] ); float err3 = ABS_ERROR(resultPtr[2], expected[2]);
float err4 = fabsf( resultPtr[3] - expected[3] ); float err4 = ABS_ERROR(resultPtr[3], expected[3]);
// Clamp to the minimum absolute error for the format // Clamp to the minimum absolute error for the format
if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; }
if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( resultPtr[0] - expected[0] ); err1 = ABS_ERROR(resultPtr[0], expected[0]);
err2 = fabsf( resultPtr[1] - expected[1] ); err2 = ABS_ERROR(resultPtr[1], expected[1]);
err3 = fabsf( resultPtr[2] - expected[2] ); err3 = ABS_ERROR(resultPtr[2], expected[2]);
err4 = fabsf( resultPtr[3] - expected[3] ); 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, &containsDenormals, lod ); imageSampler, expected, 0, &containsDenormals, lod );
float err1 = fabsf( resultPtr[0] - expected[0] ); float err1 =
float err2 = fabsf( resultPtr[1] - expected[1] ); ABS_ERROR(resultPtr[0], expected[0]);
float err3 = fabsf( resultPtr[2] - expected[2] ); float err2 =
float err4 = fabsf( resultPtr[3] - expected[3] ); 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 maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN ); float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN );
float maxErr3 = MAX( maxErr * maxPixel.p[2], 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, xOffsetValues[ j ], yOffsetValues[ j ], 0.f, norm_offset_x, norm_offset_y, 0.0f,
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( resultPtr[0] - expected[0] ); err1 = ABS_ERROR(resultPtr[0],
err2 = fabsf( resultPtr[1] - expected[1] ); expected[0]);
err3 = fabsf( resultPtr[2] - expected[2] ); err2 = ABS_ERROR(resultPtr[1],
err4 = fabsf( resultPtr[3] - expected[3] ); expected[1]);
err3 = ABS_ERROR(resultPtr[2],
expected[2]);
err4 = ABS_ERROR(resultPtr[3],
expected[3]);
} }
} }
if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) || if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) ||

122
test_conformance/images/kernel_read_write/test_read_2D_array.cpp
View File

@@ -88,8 +88,6 @@ const char *float2DArrayUnnormalizedCoordKernelSource =
static const char *samplerKernelArg = " sampler_t imageSampler,"; 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 ); extern void read_image_pixel_float( void *imageData, image_descriptor *imageInfo, int x, int y, int z, float *outData );
template <class T> int determine_validation_error_offset_2D_array( void *imagePtr, image_descriptor *imageInfo, image_sampler_data *imageSampler, template <class T> int determine_validation_error_offset_2D_array( void *imagePtr, image_descriptor *imageInfo, image_sampler_data *imageSampler,
T *resultPtr, T * expected, float error, 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, &hasDenormals, lod ); 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 // Clamp to the minimum absolute error for the format
if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; }
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN ); 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, NULL, lod ); 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, &hasDenormals, lod ); 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 ); 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 ], xOffsetValues[ j ], yOffsetValues[ j ], zOffsetValues[ j ],
imageSampler, expected, 0, NULL, lod ); 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, &hasDenormals, lod ); imageSampler, expected, 0, &hasDenormals, lod );
float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); float err1 =
float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); ABS_ERROR(sRGBmap(resultPtr[0]),
float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); sRGBmap(expected[0]));
float err4 = fabsf( resultPtr[3] - expected[3] ); 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; float maxErr = 0.5;
if( ! (err1 <= maxErr) || ! (err2 <= maxErr) || ! (err3 <= maxErr) || ! (err4 <= maxErr) ) 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); err1 =
err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); ABS_ERROR(sRGBmap(resultPtr[0]),
err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); sRGBmap(expected[0]));
err4 = fabsf( resultPtr[3] - expected[3] ); 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, &hasDenormals, lod ); imageSampler, expected, 0, &hasDenormals, lod );
float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); float err1 =
float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); ABS_ERROR(sRGBmap(resultPtr[0]),
float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); sRGBmap(expected[0]));
float err4 = fabsf( resultPtr[3] - expected[3] ); 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; float maxErr = 0.6;
if( ! (err1 <= maxErr) || ! (err2 <= maxErr) || ! (err3 <= maxErr) || ! (err4 <= maxErr) ) 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 ], xOffsetValues[ j ], yOffsetValues[ j ], zOffsetValues[ j ],
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); err1 = ABS_ERROR(
err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); sRGBmap(resultPtr[0]),
err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); sRGBmap(expected[0]));
err4 = fabsf( resultPtr[3] - expected[3] ); 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, &hasDenormals, lod ); imageSampler, expected, 0, &hasDenormals, lod );
float err1 = fabsf( resultPtr[0] - expected[0] ); float err1 =
float err2 = fabsf( resultPtr[1] - expected[1] ); ABS_ERROR(resultPtr[0], expected[0]);
float err3 = fabsf( resultPtr[2] - expected[2] ); float err2 =
float err4 = fabsf( resultPtr[3] - expected[3] ); 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 // Clamp to the minimum absolute error for the format
if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; }
if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( resultPtr[0] - expected[0] ); err1 = ABS_ERROR(resultPtr[0],
err2 = fabsf( resultPtr[1] - expected[1] ); expected[0]);
err3 = fabsf( resultPtr[2] - expected[2] ); err2 = ABS_ERROR(resultPtr[1],
err4 = fabsf( resultPtr[3] - expected[3] ); 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, &hasDenormals, lod ); imageSampler, expected, 0, &hasDenormals, lod );
float err1 = fabsf( resultPtr[0] - expected[0] ); float err1 = ABS_ERROR(resultPtr[0],
float err2 = fabsf( resultPtr[1] - expected[1] ); expected[0]);
float err3 = fabsf( resultPtr[2] - expected[2] ); float err2 = ABS_ERROR(resultPtr[1],
float err4 = fabsf( resultPtr[3] - expected[3] ); 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 maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN ); float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN );
float maxErr3 = MAX( maxErr * maxPixel.p[2], 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 ], xOffsetValues[ j ], yOffsetValues[ j ], zOffsetValues[ j ],
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( resultPtr[0] - expected[0] ); err1 = ABS_ERROR(resultPtr[0],
err2 = fabsf( resultPtr[1] - expected[1] ); expected[0]);
err3 = fabsf( resultPtr[2] - expected[2] ); err2 = ABS_ERROR(resultPtr[1],
err4 = fabsf( resultPtr[3] - expected[3] ); expected[1]);
err3 = ABS_ERROR(resultPtr[2],
expected[2]);
err4 = ABS_ERROR(resultPtr[3],
expected[3]);
} }
} }

114
test_conformance/images/kernel_read_write/test_read_3D.cpp
View File

@@ -88,8 +88,6 @@ const char *float3DUnnormalizedCoordKernelSource =
static const char *samplerKernelArg = " sampler_t imageSampler,"; 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 ); extern void read_image_pixel_float( void *imageData, image_descriptor *imageInfo, int x, int y, int z, float *outData );
template <class T> int determine_validation_error_offset( void *imagePtr, image_descriptor *imageInfo, image_sampler_data *imageSampler, template <class T> int determine_validation_error_offset( void *imagePtr, image_descriptor *imageInfo, image_sampler_data *imageSampler,
T *resultPtr, T * expected, float error, 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, &hasDenormals, lod ); imageSampler, expected, 0, &hasDenormals, lod );
float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); float err1 =
float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); ABS_ERROR(sRGBmap(resultPtr[0]),
float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); sRGBmap(expected[0]));
float err4 = fabsf( resultPtr[3] - expected[3] ); 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 // Clamp to the minimum absolute error for the format
if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; }
if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); err1 =
err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); ABS_ERROR(sRGBmap(resultPtr[0]),
err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); sRGBmap(expected[0]));
err4 = fabsf( resultPtr[3] - expected[3] ); 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, &hasDenormals, lod ); imageSampler, expected, 0, &hasDenormals, lod );
float err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); float err1 =
float err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); ABS_ERROR(sRGBmap(resultPtr[0]),
float err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); sRGBmap(expected[0]));
float err4 = fabsf( resultPtr[3] - expected[3] ); 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; float maxErr = 0.6;
if( ! (err1 <= maxErr) || ! (err2 <= maxErr) || ! (err3 <= maxErr) || ! (err4 <= maxErr) ) 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 ], xOffsetValues[ j ], yOffsetValues[ j ], zOffsetValues[ j ],
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( sRGBmap( resultPtr[0] ) - sRGBmap( expected[0] ) ); err1 = ABS_ERROR(
err2 = fabsf( sRGBmap( resultPtr[1] ) - sRGBmap( expected[1] ) ); sRGBmap(resultPtr[0]),
err3 = fabsf( sRGBmap( resultPtr[2] ) - sRGBmap( expected[2] ) ); sRGBmap(expected[0]));
err4 = fabsf( resultPtr[3] - expected[3] ); 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, &hasDenormals, lod ); imageSampler, expected, 0, &hasDenormals, lod );
float err1 = fabsf( resultPtr[0] - expected[0] ); float err1 =
float err2 = fabsf( resultPtr[1] - expected[1] ); ABS_ERROR(resultPtr[0], expected[0]);
float err3 = fabsf( resultPtr[2] - expected[2] ); float err2 =
float err4 = fabsf( resultPtr[3] - expected[3] ); 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 // Clamp to the minimum absolute error for the format
if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; } if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; }
if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( resultPtr[0] - expected[0] ); err1 = ABS_ERROR(resultPtr[0],
err2 = fabsf( resultPtr[1] - expected[1] ); expected[0]);
err3 = fabsf( resultPtr[2] - expected[2] ); err2 = ABS_ERROR(resultPtr[1],
err4 = fabsf( resultPtr[3] - expected[3] ); 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, norm_offset_x, norm_offset_y, norm_offset_z,
imageSampler, expected, 0, &hasDenormals, lod ); imageSampler, expected, 0, &hasDenormals, lod );
float err1 = fabsf( resultPtr[0] - expected[0] ); float err1 = ABS_ERROR(resultPtr[0],
float err2 = fabsf( resultPtr[1] - expected[1] ); expected[0]);
float err3 = fabsf( resultPtr[2] - expected[2] ); float err2 = ABS_ERROR(resultPtr[1],
float err4 = fabsf( resultPtr[3] - expected[3] ); 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 maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN ); float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN );
float maxErr3 = MAX( maxErr * maxPixel.p[2], 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 ], xOffsetValues[ j ], yOffsetValues[ j ], zOffsetValues[ j ],
imageSampler, expected, 0, NULL, lod ); imageSampler, expected, 0, NULL, lod );
err1 = fabsf( resultPtr[0] - expected[0] ); err1 = ABS_ERROR(resultPtr[0],
err2 = fabsf( resultPtr[1] - expected[1] ); expected[0]);
err3 = fabsf( resultPtr[2] - expected[2] ); err2 = ABS_ERROR(resultPtr[1],
err4 = fabsf( resultPtr[3] - expected[3] ); 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; return 0;
} }

121
test_conformance/math_brute_force/main.cpp
View File

@@ -15,10 +15,9 @@
// //
#include "Utility.h" #include "Utility.h"
#include <stdio.h> #include <cstdio>
#include <string.h> #include <cstdlib>
#include <string>
#include <stdlib.h>
#include <time.h> #include <time.h>
#include "FunctionList.h" #include "FunctionList.h"
#include "Sleep.h" #include "Sleep.h"
@@ -106,24 +105,6 @@ cl_device_fp_config gDoubleCapabilities = 0;
int gWimpyReductionFactor = 32; int gWimpyReductionFactor = 32;
int gWimpyBufferSize = BUFFER_SIZE; int gWimpyBufferSize = BUFFER_SIZE;
int gVerboseBruteForce = 0; 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 int ParseArgs( int argc, const char **argv );
static void PrintUsage( void ); static void PrintUsage( void );
@@ -255,44 +236,6 @@ int doTest( const char* name )
//Re-enable testing fast-relaxed-math mode //Re-enable testing fast-relaxed-math mode
gTestFastRelaxed = testFastRelaxedTmp; 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; return error;
@@ -1179,20 +1122,20 @@ test_status InitCL( cl_device_id device )
if( DOUBLE_REQUIRED_FEATURES != (gDoubleCapabilities & DOUBLE_REQUIRED_FEATURES) ) if( DOUBLE_REQUIRED_FEATURES != (gDoubleCapabilities & DOUBLE_REQUIRED_FEATURES) )
{ {
char list[300] = ""; std::string list;
if( 0 == (gDoubleCapabilities & CL_FP_FMA) ) if (0 == (gDoubleCapabilities & CL_FP_FMA)) list += "CL_FP_FMA, ";
strncat( list, "CL_FP_FMA, ", sizeof( list )-1 );
if( 0 == (gDoubleCapabilities & CL_FP_ROUND_TO_NEAREST) ) 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) ) 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) ) 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) ) 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) ) if( 0 == (gDoubleCapabilities & CL_FP_DENORM) )
strncat( list, "CL_FP_DENORM, ", sizeof( list )-1 ); list += "CL_FP_DENORM, ";
vlog_error( "ERROR: required double features are missing: %s\n", list ); vlog_error("ERROR: required double features are missing: %s\n",
list.c_str());
return TEST_FAIL; return TEST_FAIL;
} }
@@ -1201,43 +1144,6 @@ test_status InitCL( cl_device_id device )
return TEST_FAIL; return TEST_FAIL;
#endif #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 ); configSize = sizeof( gDeviceFrequency );
if( (error = clGetDeviceInfo( gDevice, CL_DEVICE_MAX_CLOCK_FREQUENCY, configSize, &gDeviceFrequency, NULL )) ) 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 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" ); 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] ); vlog( "\tIs Embedded? %s\n", no_yes[0 != gIsEmbedded] );
if( gIsEmbedded ) if( gIsEmbedded )

33
test_conformance/math_brute_force/reference_math.cpp
View File

@@ -37,39 +37,6 @@
#define EVALUATE( x ) x #define EVALUATE( x ) x
#define CONCATENATE(x, y) x ## EVALUATE(y) #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 #pragma STDC FP_CONTRACT OFF
static void __log2_ep(double *hi, double *lo, double x); static void __log2_ep(double *hi, double *lo, double x);

14
test_conformance/spir/datagen.h
View File

@@ -274,20 +274,16 @@ public:
if (CL_SUCCESS != error) if (CL_SUCCESS != error)
throw Exceptions::TestError("clGetSupportedImageFormats failed\n", error); throw Exceptions::TestError("clGetSupportedImageFormats failed\n", error);
std::auto_ptr<cl_image_format> supportedFormats(new cl_image_format[actualNumFormats]); std::vector<cl_image_format> supportedFormats(actualNumFormats);
error = clGetSupportedImageFormats( error = clGetSupportedImageFormats(context, flags, m_desc.image_type,
context, actualNumFormats,
flags, supportedFormats.data(), NULL);
m_desc.image_type,
actualNumFormats,
supportedFormats.get(),
NULL);
if (CL_SUCCESS != error) if (CL_SUCCESS != error)
throw Exceptions::TestError("clGetSupportedImageFormats failed\n", error); throw Exceptions::TestError("clGetSupportedImageFormats failed\n", error);
for (size_t i=0; i<actualNumFormats; ++i) for (size_t i=0; i<actualNumFormats; ++i)
{ {
cl_image_format curFormat = supportedFormats.get()[i]; cl_image_format curFormat = supportedFormats[i];
if(imgFormat.image_channel_order == curFormat.image_channel_order && if(imgFormat.image_channel_order == curFormat.image_channel_order &&
imgFormat.image_channel_data_type == curFormat.image_channel_data_type) imgFormat.image_channel_data_type == curFormat.image_channel_data_type)

4
test_conformance/spir/main.cpp
View File

@@ -147,8 +147,8 @@ static void get_spir_version(cl_device_id device, std::vector<float>& versions)
cl_int err; cl_int err;
size_t size = 0; size_t size = 0;
if (err = clGetDeviceInfo(device, CL_DEVICE_SPIR_VERSIONS, sizeof(version), if ((err = clGetDeviceInfo(device, CL_DEVICE_SPIR_VERSIONS, sizeof(version),
(void*)version, &size)) (void *)version, &size)))
{ {
log_error( "Error: failed to obtain SPIR version at %s:%d (err = %d)\n", log_error( "Error: failed to obtain SPIR version at %s:%d (err = %d)\n",
__FILE__, __LINE__, err ); __FILE__, __LINE__, err );

2
test_conformance/spir/run_build_test.cpp
View File

@@ -268,7 +268,7 @@ static bool run_test(cl_context context, cl_command_queue queue, cl_program clpr
{ {
WorkSizeInfo ws; WorkSizeInfo ws;
TestResult cl_result; TestResult cl_result;
std::auto_ptr<TestResult> bc_result; std::unique_ptr<TestResult> bc_result;
// first, run the single CL test // first, run the single CL test
{ {
// make sure that the kernel will be released before the program // make sure that the kernel will be released before the program