diff --git a/test_conformance/CMakeLists.txt b/test_conformance/CMakeLists.txt index e488b705..e93d83e2 100644 --- a/test_conformance/CMakeLists.txt +++ b/test_conformance/CMakeLists.txt @@ -9,7 +9,6 @@ add_subdirectory( atomics ) add_subdirectory( basic ) add_subdirectory( buffers ) add_subdirectory( commonfns ) -add_subdirectory( compatibility ) add_subdirectory( compiler ) add_subdirectory( computeinfo ) add_subdirectory( contractions ) diff --git a/test_conformance/compatibility/CMakeLists.txt b/test_conformance/compatibility/CMakeLists.txt deleted file mode 100644 index ad621bfa..00000000 --- a/test_conformance/compatibility/CMakeLists.txt +++ /dev/null @@ -1,2 +0,0 @@ -add_subdirectory(test_common) -add_subdirectory(test_conformance) diff --git a/test_conformance/compatibility/test_common/CMakeLists.txt b/test_conformance/compatibility/test_common/CMakeLists.txt deleted file mode 100644 index 0a3f5ded..00000000 --- a/test_conformance/compatibility/test_common/CMakeLists.txt +++ /dev/null @@ -1,14 +0,0 @@ - -set(HARNESS_COMPAT_SOURCES - harness/errorHelpers.c - harness/imageHelpers.cpp - harness/kernelHelpers.c - harness/testHarness.c -) - -set_source_files_properties(${HARNESS_COMPAT_SOURCES} PROPERTIES LANGUAGE CXX) - -add_library(harness-compat STATIC ${HARNESS_COMPAT_SOURCES}) - -target_link_libraries(harness-compat harness-common) - diff --git a/test_conformance/compatibility/test_common/harness/clImageHelper.h b/test_conformance/compatibility/test_common/harness/clImageHelper.h deleted file mode 100644 index 8537ddcd..00000000 --- a/test_conformance/compatibility/test_common/harness/clImageHelper.h +++ /dev/null @@ -1,253 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#ifndef test_conformance_clImageHelper_h -#define test_conformance_clImageHelper_h - -#ifdef __APPLE__ -#include -#else -#include -#endif - -#include -#include "errorHelpers.h" - -#ifdef __cplusplus -extern "C" { -#endif - - - // helper function to replace clCreateImage2D , to make the existing code use - // the functions of version 1.2 and veriosn 1.1 respectively - - inline cl_mem create_image_2d (cl_context context, - cl_mem_flags flags, - const cl_image_format *image_format, - size_t image_width, - size_t image_height, - size_t image_row_pitch, - void *host_ptr, - cl_int *errcode_ret) - { - cl_mem mImage = NULL; - -#ifdef CL_VERSION_1_2 - cl_image_desc image_desc_dest; - image_desc_dest.image_type = CL_MEM_OBJECT_IMAGE2D;; - image_desc_dest.image_width = image_width; - image_desc_dest.image_height = image_height; - image_desc_dest.image_depth= 0;// not usedfor 2d - image_desc_dest.image_array_size = 0;// not used for 2d - image_desc_dest.image_row_pitch = image_row_pitch; - image_desc_dest.image_slice_pitch = 0; - image_desc_dest.num_mip_levels = 0; - image_desc_dest.num_samples = 0; - image_desc_dest.buffer = NULL;// no image type of CL_MEM_OBJECT_IMAGE1D_BUFFER in CL_VERSION_1_1, so always is NULL - mImage = clCreateImage( context, flags, image_format, &image_desc_dest, host_ptr, errcode_ret ); - if (errcode_ret && (*errcode_ret)) { - // Log an info message and rely on the calling function to produce an error - // if necessary. - log_info("clCreateImage failed (%d)\n", *errcode_ret); - } - -#else - mImage = clCreateImage2D( context, flags, image_format, image_width, image_height, image_row_pitch, host_ptr, errcode_ret ); - if (errcode_ret && (*errcode_ret)) { - // Log an info message and rely on the calling function to produce an error - // if necessary. - log_info("clCreateImage2D failed (%d)\n", *errcode_ret); - } -#endif - - return mImage; - } - - inline cl_mem create_image_3d (cl_context context, - cl_mem_flags flags, - const cl_image_format *image_format, - size_t image_width, - size_t image_height, - size_t image_depth, - size_t image_row_pitch, - size_t image_slice_pitch, - void *host_ptr, - cl_int *errcode_ret) - { - cl_mem mImage; - -#ifdef CL_VERSION_1_2 - cl_image_desc image_desc; - image_desc.image_type = CL_MEM_OBJECT_IMAGE3D; - image_desc.image_width = image_width; - image_desc.image_height = image_height; - image_desc.image_depth = image_depth; - image_desc.image_array_size = 0;// not used for one image - image_desc.image_row_pitch = image_row_pitch; - image_desc.image_slice_pitch = image_slice_pitch; - image_desc.num_mip_levels = 0; - image_desc.num_samples = 0; - image_desc.buffer = NULL; // no image type of CL_MEM_OBJECT_IMAGE1D_BUFFER in CL_VERSION_1_1, so always is NULL - mImage = clCreateImage( context, - flags, - image_format, - &image_desc, - host_ptr, - errcode_ret ); - if (errcode_ret && (*errcode_ret)) { - // Log an info message and rely on the calling function to produce an error - // if necessary. - log_info("clCreateImage failed (%d)\n", *errcode_ret); - } - -#else - mImage = clCreateImage3D( context, - flags, image_format, - image_width, - image_height, - image_depth, - image_row_pitch, - image_slice_pitch, - host_ptr, - errcode_ret ); - if (errcode_ret && (*errcode_ret)) { - // Log an info message and rely on the calling function to produce an error - // if necessary. - log_info("clCreateImage3D failed (%d)\n", *errcode_ret); - } -#endif - - return mImage; - } - - inline cl_mem create_image_2d_array (cl_context context, - cl_mem_flags flags, - const cl_image_format *image_format, - size_t image_width, - size_t image_height, - size_t image_array_size, - size_t image_row_pitch, - size_t image_slice_pitch, - void *host_ptr, - cl_int *errcode_ret) - { - cl_mem mImage; - - cl_image_desc image_desc; - image_desc.image_type = CL_MEM_OBJECT_IMAGE2D_ARRAY; - image_desc.image_width = image_width; - image_desc.image_height = image_height; - image_desc.image_depth = 1; - image_desc.image_array_size = image_array_size; - image_desc.image_row_pitch = image_row_pitch; - image_desc.image_slice_pitch = image_slice_pitch; - image_desc.num_mip_levels = 0; - image_desc.num_samples = 0; - image_desc.buffer = NULL; - mImage = clCreateImage( context, - flags, - image_format, - &image_desc, - host_ptr, - errcode_ret ); - if (errcode_ret && (*errcode_ret)) { - // Log an info message and rely on the calling function to produce an error - // if necessary. - log_info("clCreateImage failed (%d)\n", *errcode_ret); - } - - return mImage; - } - - inline cl_mem create_image_1d_array (cl_context context, - cl_mem_flags flags, - const cl_image_format *image_format, - size_t image_width, - size_t image_array_size, - size_t image_row_pitch, - size_t image_slice_pitch, - void *host_ptr, - cl_int *errcode_ret) - { - cl_mem mImage; - - cl_image_desc image_desc; - image_desc.image_type = CL_MEM_OBJECT_IMAGE1D_ARRAY; - image_desc.image_width = image_width; - image_desc.image_height = 1; - image_desc.image_depth = 1; - image_desc.image_array_size = image_array_size; - image_desc.image_row_pitch = image_row_pitch; - image_desc.image_slice_pitch = image_slice_pitch; - image_desc.num_mip_levels = 0; - image_desc.num_samples = 0; - image_desc.buffer = NULL; - mImage = clCreateImage( context, - flags, - image_format, - &image_desc, - host_ptr, - errcode_ret ); - if (errcode_ret && (*errcode_ret)) { - // Log an info message and rely on the calling function to produce an error - // if necessary. - log_info("clCreateImage failed (%d)\n", *errcode_ret); - } - - return mImage; - } - - inline cl_mem create_image_1d (cl_context context, - cl_mem_flags flags, - const cl_image_format *image_format, - size_t image_width, - size_t image_row_pitch, - void *host_ptr, - cl_mem buffer, - cl_int *errcode_ret) - { - cl_mem mImage; - - cl_image_desc image_desc; - image_desc.image_type = buffer ? CL_MEM_OBJECT_IMAGE1D_BUFFER: CL_MEM_OBJECT_IMAGE1D; - image_desc.image_width = image_width; - image_desc.image_height = 1; - image_desc.image_depth = 1; - image_desc.image_row_pitch = image_row_pitch; - image_desc.image_slice_pitch = 0; - image_desc.num_mip_levels = 0; - image_desc.num_samples = 0; - image_desc.buffer = buffer; - mImage = clCreateImage( context, - flags, - image_format, - &image_desc, - host_ptr, - errcode_ret ); - if (errcode_ret && (*errcode_ret)) { - // Log an info message and rely on the calling function to produce an error - // if necessary. - log_info("clCreateImage failed (%d)\n", *errcode_ret); - } - - return mImage; - } - - -#ifdef __cplusplus -} -#endif - -#endif diff --git a/test_conformance/compatibility/test_common/harness/errorHelpers.c b/test_conformance/compatibility/test_common/harness/errorHelpers.c deleted file mode 100644 index 79c4ca21..00000000 --- a/test_conformance/compatibility/test_common/harness/errorHelpers.c +++ /dev/null @@ -1,579 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "compat.h" - -#include -#include -#include - -#include "errorHelpers.h" - -const char *IGetErrorString( int clErrorCode ) -{ - switch( clErrorCode ) - { - case CL_SUCCESS: return "CL_SUCCESS"; - case CL_DEVICE_NOT_FOUND: return "CL_DEVICE_NOT_FOUND"; - case CL_DEVICE_NOT_AVAILABLE: return "CL_DEVICE_NOT_AVAILABLE"; - case CL_COMPILER_NOT_AVAILABLE: return "CL_COMPILER_NOT_AVAILABLE"; - case CL_MEM_OBJECT_ALLOCATION_FAILURE: return "CL_MEM_OBJECT_ALLOCATION_FAILURE"; - case CL_OUT_OF_RESOURCES: return "CL_OUT_OF_RESOURCES"; - case CL_OUT_OF_HOST_MEMORY: return "CL_OUT_OF_HOST_MEMORY"; - case CL_PROFILING_INFO_NOT_AVAILABLE: return "CL_PROFILING_INFO_NOT_AVAILABLE"; - case CL_MEM_COPY_OVERLAP: return "CL_MEM_COPY_OVERLAP"; - case CL_IMAGE_FORMAT_MISMATCH: return "CL_IMAGE_FORMAT_MISMATCH"; - case CL_IMAGE_FORMAT_NOT_SUPPORTED: return "CL_IMAGE_FORMAT_NOT_SUPPORTED"; - case CL_BUILD_PROGRAM_FAILURE: return "CL_BUILD_PROGRAM_FAILURE"; - case CL_MAP_FAILURE: return "CL_MAP_FAILURE"; - case CL_MISALIGNED_SUB_BUFFER_OFFSET: return "CL_MISALIGNED_SUB_BUFFER_OFFSET"; - case CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST: return "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST"; - case CL_COMPILE_PROGRAM_FAILURE: return "CL_COMPILE_PROGRAM_FAILURE"; - case CL_LINKER_NOT_AVAILABLE: return "CL_LINKER_NOT_AVAILABLE"; - case CL_LINK_PROGRAM_FAILURE: return "CL_LINK_PROGRAM_FAILURE"; - case CL_DEVICE_PARTITION_FAILED: return "CL_DEVICE_PARTITION_FAILED"; - case CL_KERNEL_ARG_INFO_NOT_AVAILABLE: return "CL_KERNEL_ARG_INFO_NOT_AVAILABLE"; - case CL_INVALID_VALUE: return "CL_INVALID_VALUE"; - case CL_INVALID_DEVICE_TYPE: return "CL_INVALID_DEVICE_TYPE"; - case CL_INVALID_DEVICE: return "CL_INVALID_DEVICE"; - case CL_INVALID_CONTEXT: return "CL_INVALID_CONTEXT"; - case CL_INVALID_QUEUE_PROPERTIES: return "CL_INVALID_QUEUE_PROPERTIES"; - case CL_INVALID_COMMAND_QUEUE: return "CL_INVALID_COMMAND_QUEUE"; - case CL_INVALID_HOST_PTR: return "CL_INVALID_HOST_PTR"; - case CL_INVALID_MEM_OBJECT: return "CL_INVALID_MEM_OBJECT"; - case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR: return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"; - case CL_INVALID_IMAGE_SIZE: return "CL_INVALID_IMAGE_SIZE"; - case CL_INVALID_SAMPLER: return "CL_INVALID_SAMPLER"; - case CL_INVALID_BINARY: return "CL_INVALID_BINARY"; - case CL_INVALID_BUILD_OPTIONS: return "CL_INVALID_BUILD_OPTIONS"; - case CL_INVALID_PROGRAM: return "CL_INVALID_PROGRAM"; - case CL_INVALID_PROGRAM_EXECUTABLE: return "CL_INVALID_PROGRAM_EXECUTABLE"; - case CL_INVALID_KERNEL_NAME: return "CL_INVALID_KERNEL_NAME"; - case CL_INVALID_KERNEL_DEFINITION: return "CL_INVALID_KERNEL_DEFINITION"; - case CL_INVALID_KERNEL: return "CL_INVALID_KERNEL"; - case CL_INVALID_ARG_INDEX: return "CL_INVALID_ARG_INDEX"; - case CL_INVALID_ARG_VALUE: return "CL_INVALID_ARG_VALUE"; - case CL_INVALID_ARG_SIZE: return "CL_INVALID_ARG_SIZE"; - case CL_INVALID_KERNEL_ARGS: return "CL_INVALID_KERNEL_ARGS"; - case CL_INVALID_WORK_DIMENSION: return "CL_INVALID_WORK_DIMENSION"; - case CL_INVALID_WORK_GROUP_SIZE: return "CL_INVALID_WORK_GROUP_SIZE"; - case CL_INVALID_WORK_ITEM_SIZE: return "CL_INVALID_WORK_ITEM_SIZE"; - case CL_INVALID_GLOBAL_OFFSET: return "CL_INVALID_GLOBAL_OFFSET"; - case CL_INVALID_EVENT_WAIT_LIST: return "CL_INVALID_EVENT_WAIT_LIST"; - case CL_INVALID_EVENT: return "CL_INVALID_EVENT"; - case CL_INVALID_OPERATION: return "CL_INVALID_OPERATION"; - case CL_INVALID_GL_OBJECT: return "CL_INVALID_GL_OBJECT"; - case CL_INVALID_BUFFER_SIZE: return "CL_INVALID_BUFFER_SIZE"; - case CL_INVALID_MIP_LEVEL: return "CL_INVALID_MIP_LEVEL"; - case CL_INVALID_GLOBAL_WORK_SIZE: return "CL_INVALID_GLOBAL_WORK_SIZE"; - case CL_INVALID_PROPERTY: return "CL_INVALID_PROPERTY"; - case CL_INVALID_IMAGE_DESCRIPTOR: return "CL_INVALID_IMAGE_DESCRIPTOR"; - case CL_INVALID_COMPILER_OPTIONS: return "CL_INVALID_COMPILER_OPTIONS"; - case CL_INVALID_LINKER_OPTIONS: return "CL_INVALID_LINKER_OPTIONS"; - case CL_INVALID_DEVICE_PARTITION_COUNT: return "CL_INVALID_DEVICE_PARTITION_COUNT"; - default: return "(unknown)"; - } -} - -const char *GetChannelOrderName( cl_channel_order order ) -{ - switch( order ) - { - case CL_R: return "CL_R"; - case CL_A: return "CL_A"; - case CL_Rx: return "CL_Rx"; - case CL_RG: return "CL_RG"; - case CL_RA: return "CL_RA"; - case CL_RGx: return "CL_RGx"; - case CL_RGB: return "CL_RGB"; - case CL_RGBx: return "CL_RGBx"; - case CL_RGBA: return "CL_RGBA"; - case CL_ARGB: return "CL_ARGB"; - case CL_BGRA: return "CL_BGRA"; - case CL_INTENSITY: return "CL_INTENSITY"; - case CL_LUMINANCE: return "CL_LUMINANCE"; -#if defined CL_1RGB_APPLE - case CL_1RGB_APPLE: return "CL_1RGB_APPLE"; -#endif -#if defined CL_BGR1_APPLE - case CL_BGR1_APPLE: return "CL_BGR1_APPLE"; -#endif - default: return NULL; - } -} - -int IsChannelOrderSupported( cl_channel_order order ) -{ - switch( order ) - { - case CL_R: - case CL_A: - case CL_Rx: - case CL_RG: - case CL_RA: - case CL_RGx: - case CL_RGB: - case CL_RGBx: - case CL_RGBA: - case CL_ARGB: - case CL_BGRA: - case CL_INTENSITY: - case CL_LUMINANCE: - return 1; -#if defined CL_1RGB_APPLE - case CL_1RGB_APPLE: - return 1; -#endif -#if defined CL_BGR1_APPLE - case CL_BGR1_APPLE: - return 1; -#endif - default: - return 0; - } -} - -const char *GetChannelTypeName( cl_channel_type type ) -{ - switch( type ) - { - case CL_SNORM_INT8: return "CL_SNORM_INT8"; - case CL_SNORM_INT16: return "CL_SNORM_INT16"; - case CL_UNORM_INT8: return "CL_UNORM_INT8"; - case CL_UNORM_INT16: return "CL_UNORM_INT16"; - case CL_UNORM_SHORT_565: return "CL_UNORM_SHORT_565"; - case CL_UNORM_SHORT_555: return "CL_UNORM_SHORT_555"; - case CL_UNORM_INT_101010: return "CL_UNORM_INT_101010"; - case CL_SIGNED_INT8: return "CL_SIGNED_INT8"; - case CL_SIGNED_INT16: return "CL_SIGNED_INT16"; - case CL_SIGNED_INT32: return "CL_SIGNED_INT32"; - case CL_UNSIGNED_INT8: return "CL_UNSIGNED_INT8"; - case CL_UNSIGNED_INT16: return "CL_UNSIGNED_INT16"; - case CL_UNSIGNED_INT32: return "CL_UNSIGNED_INT32"; - case CL_HALF_FLOAT: return "CL_HALF_FLOAT"; - case CL_FLOAT: return "CL_FLOAT"; -#ifdef CL_SFIXED14_APPLE - case CL_SFIXED14_APPLE: return "CL_SFIXED14_APPLE"; -#endif - default: return NULL; - } -} - -int IsChannelTypeSupported( cl_channel_type type ) -{ - switch( type ) - { - case CL_SNORM_INT8: - case CL_SNORM_INT16: - case CL_UNORM_INT8: - case CL_UNORM_INT16: - case CL_UNORM_SHORT_565: - case CL_UNORM_SHORT_555: - case CL_UNORM_INT_101010: - case CL_SIGNED_INT8: - case CL_SIGNED_INT16: - case CL_SIGNED_INT32: - case CL_UNSIGNED_INT8: - case CL_UNSIGNED_INT16: - case CL_UNSIGNED_INT32: - case CL_HALF_FLOAT: - case CL_FLOAT: - return 1; -#ifdef CL_SFIXED14_APPLE - case CL_SFIXED14_APPLE: - return 1; -#endif - default: - return 0; - } -} - -const char *GetAddressModeName( cl_addressing_mode mode ) -{ - switch( mode ) - { - case CL_ADDRESS_NONE: return "CL_ADDRESS_NONE"; - case CL_ADDRESS_CLAMP_TO_EDGE: return "CL_ADDRESS_CLAMP_TO_EDGE"; - case CL_ADDRESS_CLAMP: return "CL_ADDRESS_CLAMP"; - case CL_ADDRESS_REPEAT: return "CL_ADDRESS_REPEAT"; - case CL_ADDRESS_MIRRORED_REPEAT: return "CL_ADDRESS_MIRRORED_REPEAT"; - default: return NULL; - } -} - -const char *GetDeviceTypeName( cl_device_type type ) -{ - switch( type ) - { - case CL_DEVICE_TYPE_GPU: return "CL_DEVICE_TYPE_GPU"; - case CL_DEVICE_TYPE_CPU: return "CL_DEVICE_TYPE_CPU"; - case CL_DEVICE_TYPE_ACCELERATOR: return "CL_DEVICE_TYPE_ACCELERATOR"; - case CL_DEVICE_TYPE_ALL: return "CL_DEVICE_TYPE_ALL"; - default: return NULL; - } -} - -const char *GetDataVectorString( void *dataBuffer, size_t typeSize, size_t vecSize, char *buffer ) -{ - static char scratch[ 1024 ]; - size_t i, j; - - if( buffer == NULL ) - buffer = scratch; - - unsigned char *p = (unsigned char *)dataBuffer; - char *bPtr; - - buffer[ 0 ] = 0; - bPtr = buffer; - for( i = 0; i < vecSize; i++ ) - { - if( i > 0 ) - { - bPtr[ 0 ] = ' '; - bPtr++; - } - for( j = 0; j < typeSize; j++ ) - { - sprintf( bPtr, "%02x", (unsigned int)p[ typeSize - j - 1 ] ); - bPtr += 2; - } - p += typeSize; - } - bPtr[ 0 ] = 0; - - return buffer; -} - -#ifndef MAX -#define MAX( _a, _b ) ((_a) > (_b) ? (_a) : (_b)) -#endif - -#if defined( _MSC_VER ) -#define scalbnf(_a, _i ) ldexpf( _a, _i ) -#define scalbn(_a, _i ) ldexp( _a, _i ) -#define scalbnl(_a, _i ) ldexpl( _a, _i ) -#endif - -static float Ulp_Error_Half_Float( float test, double reference ); -static inline float half2float( cl_ushort half ); - -// taken from math tests -#define HALF_MIN_EXP -13 -#define HALF_MANT_DIG 11 -static float Ulp_Error_Half_Float( float test, double reference ) -{ - union{ double d; uint64_t u; }u; u.d = reference; - - // Note: This function presumes that someone has already tested whether the result is correctly, - // rounded before calling this function. That test: - // - // if( (float) reference == test ) - // return 0.0f; - // - // would ensure that cases like fabs(reference) > FLT_MAX are weeded out before we get here. - // Otherwise, we'll return inf ulp error here, for what are otherwise correctly rounded - // results. - - double testVal = test; - if( u.u & 0x000fffffffffffffULL ) - { // Non-power of two and NaN - if( isnan( reference ) && isnan( test ) ) - return 0.0f; // if we are expecting a NaN, any NaN is fine - - // The unbiased exponent of the ulp unit place - int ulp_exp = HALF_MANT_DIG - 1 - MAX( ilogb( reference), HALF_MIN_EXP-1 ); - - // Scale the exponent of the error - return (float) scalbn( testVal - reference, ulp_exp ); - } - - if( isinf( reference ) ) - { - if( (double) test == reference ) - return 0.0f; - - return (float) (testVal - reference ); - } - - // reference is a normal power of two or a zero - int ulp_exp = HALF_MANT_DIG - 1 - MAX( ilogb( reference) - 1, HALF_MIN_EXP-1 ); - - // Scale the exponent of the error - return (float) scalbn( testVal - reference, ulp_exp ); -} - -// Taken from vLoadHalf test -static inline float half2float( cl_ushort us ) -{ - uint32_t u = us; - uint32_t sign = (u << 16) & 0x80000000; - int32_t exponent = (u & 0x7c00) >> 10; - uint32_t mantissa = (u & 0x03ff) << 13; - union{ unsigned int u; float f;}uu; - - if( exponent == 0 ) - { - if( mantissa == 0 ) - return sign ? -0.0f : 0.0f; - - int shift = __builtin_clz( mantissa ) - 8; - exponent -= shift-1; - mantissa <<= shift; - mantissa &= 0x007fffff; - } - else - if( exponent == 31) - { - uu.u = mantissa | sign; - if( mantissa ) - uu.u |= 0x7fc00000; - else - uu.u |= 0x7f800000; - - return uu.f; - } - - exponent += 127 - 15; - exponent <<= 23; - - exponent |= mantissa; - uu.u = exponent | sign; - - return uu.f; -} - -float Ulp_Error_Half( cl_ushort test, float reference ) -{ - return Ulp_Error_Half_Float( half2float(test), reference ); -} - - -float Ulp_Error( float test, double reference ) -{ - union{ double d; uint64_t u; }u; u.d = reference; - double testVal = test; - - // Note: This function presumes that someone has already tested whether the result is correctly, - // rounded before calling this function. That test: - // - // if( (float) reference == test ) - // return 0.0f; - // - // would ensure that cases like fabs(reference) > FLT_MAX are weeded out before we get here. - // Otherwise, we'll return inf ulp error here, for what are otherwise correctly rounded - // results. - - - if( isinf( reference ) ) - { - if( testVal == reference ) - return 0.0f; - - return (float) (testVal - reference ); - } - - if( isinf( testVal) ) - { // infinite test value, but finite (but possibly overflowing in float) reference. - // - // The function probably overflowed prematurely here. Formally, the spec says this is - // an infinite ulp error and should not be tolerated. Unfortunately, this would mean - // that the internal precision of some half_pow implementations would have to be 29+ bits - // at half_powr( 0x1.fffffep+31, 4) to correctly determine that 4*log2( 0x1.fffffep+31 ) - // is not exactly 128.0. You might represent this for example as 4*(32 - ~2**-24), which - // after rounding to single is 4*32 = 128, which will ultimately result in premature - // overflow, even though a good faith representation would be correct to within 2**-29 - // interally. - - // In the interest of not requiring the implementation go to extraordinary lengths to - // deliver a half precision function, we allow premature overflow within the limit - // of the allowed ulp error. Towards, that end, we "pretend" the test value is actually - // 2**128, the next value that would appear in the number line if float had sufficient range. - testVal = copysign( MAKE_HEX_DOUBLE(0x1.0p128, 0x1LL, 128), testVal ); - - // Note that the same hack may not work in long double, which is not guaranteed to have - // more range than double. It is not clear that premature overflow should be tolerated for - // double. - } - - if( u.u & 0x000fffffffffffffULL ) - { // Non-power of two and NaN - if( isnan( reference ) && isnan( test ) ) - return 0.0f; // if we are expecting a NaN, any NaN is fine - - // The unbiased exponent of the ulp unit place - int ulp_exp = FLT_MANT_DIG - 1 - MAX( ilogb( reference), FLT_MIN_EXP-1 ); - - // Scale the exponent of the error - return (float) scalbn( testVal - reference, ulp_exp ); - } - - // reference is a normal power of two or a zero - // The unbiased exponent of the ulp unit place - int ulp_exp = FLT_MANT_DIG - 1 - MAX( ilogb( reference) - 1, FLT_MIN_EXP-1 ); - - // Scale the exponent of the error - return (float) scalbn( testVal - reference, ulp_exp ); -} - -float Ulp_Error_Double( double test, long double reference ) -{ - // Deal with long double = double - // On most systems long double is a higher precision type than double. They provide either - // a 80-bit or greater floating point type, or they provide a head-tail double double format. - // That is sufficient to represent the accuracy of a floating point result to many more bits - // than double and we can calculate sub-ulp errors. This is the standard system for which this - // test suite is designed. - // - // On some systems double and long double are the same thing. Then we run into a problem, - // because our representation of the infinitely precise result (passed in as reference above) - // can be off by as much as a half double precision ulp itself. In this case, we inflate the - // reported error by half an ulp to take this into account. A more correct and permanent fix - // would be to undertake refactoring the reference code to return results in this format: - // - // typedef struct DoubleReference - // { // true value = correctlyRoundedResult + ulps * ulp(correctlyRoundedResult) (infinitely precise) - // double correctlyRoundedResult; // as best we can - // double ulps; // plus a fractional amount to account for the difference - // }DoubleReference; // between infinitely precise result and correctlyRoundedResult, in units of ulps. - // - // This would provide a useful higher-than-double precision format for everyone that we can use, - // and would solve a few problems with representing absolute errors below DBL_MIN and over DBL_MAX for systems - // that use a head to tail double double for long double. - - // Note: This function presumes that someone has already tested whether the result is correctly, - // rounded before calling this function. That test: - // - // if( (float) reference == test ) - // return 0.0f; - // - // would ensure that cases like fabs(reference) > FLT_MAX are weeded out before we get here. - // Otherwise, we'll return inf ulp error here, for what are otherwise correctly rounded - // results. - - - int x; - long double testVal = test; - if( 0.5L != frexpl( reference, &x) ) - { // Non-power of two and NaN - if( isinf( reference ) ) - { - if( testVal == reference ) - return 0.0f; - - return (float) ( testVal - reference ); - } - - if( isnan( reference ) && isnan( test ) ) - return 0.0f; // if we are expecting a NaN, any NaN is fine - - // The unbiased exponent of the ulp unit place - int ulp_exp = DBL_MANT_DIG - 1 - MAX( ilogbl( reference), DBL_MIN_EXP-1 ); - - // Scale the exponent of the error - float result = (float) scalbnl( testVal - reference, ulp_exp ); - - // account for rounding error in reference result on systems that do not have a higher precision floating point type (see above) - if( sizeof(long double) == sizeof( double ) ) - result += copysignf( 0.5f, result); - - return result; - - } - - // reference is a normal power of two or a zero - // The unbiased exponent of the ulp unit place - int ulp_exp = DBL_MANT_DIG - 1 - MAX( ilogbl( reference) - 1, DBL_MIN_EXP-1 ); - - // Scale the exponent of the error - float result = (float) scalbnl( testVal - reference, ulp_exp ); - - // account for rounding error in reference result on systems that do not have a higher precision floating point type (see above) - if( sizeof(long double) == sizeof( double ) ) - result += copysignf( 0.5f, result); - - return result; -} - -cl_int OutputBuildLogs(cl_program program, cl_uint num_devices, cl_device_id *device_list) -{ - int error; - size_t size_ret; - - // Does the program object exist? - if (program != NULL) { - - // Was the number of devices given - if (num_devices == 0) { - - // If zero devices were specified then allocate and query the device list from the context - cl_context context; - error = clGetProgramInfo(program, CL_PROGRAM_CONTEXT, sizeof(context), &context, NULL); - test_error( error, "Unable to query program's context" ); - error = clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &size_ret); - test_error( error, "Unable to query context's device size" ); - num_devices = size_ret / sizeof(cl_device_id); - device_list = (cl_device_id *) malloc(size_ret); - if (device_list == NULL) { - print_error( error, "malloc failed" ); - return CL_OUT_OF_HOST_MEMORY; - } - error = clGetContextInfo(context, CL_CONTEXT_DEVICES, size_ret, device_list, NULL); - test_error( error, "Unable to query context's devices" ); - - } - - // For each device in the device_list - unsigned int i; - for (i = 0; i < num_devices; i++) { - - // Get the build status - cl_build_status build_status; - error = clGetProgramBuildInfo(program, - device_list[i], - CL_PROGRAM_BUILD_STATUS, - sizeof(build_status), - &build_status, - &size_ret); - test_error( error, "Unable to query build status" ); - - // If the build failed then log the status, and allocate the build log, log it and free it - if (build_status != CL_BUILD_SUCCESS) { - - log_error("ERROR: CL_PROGRAM_BUILD_STATUS=%d\n", (int) build_status); - error = clGetProgramBuildInfo(program, device_list[i], CL_PROGRAM_BUILD_LOG, 0, NULL, &size_ret); - test_error( error, "Unable to query build log size" ); - char *build_log = (char *) malloc(size_ret); - error = clGetProgramBuildInfo(program, device_list[i], CL_PROGRAM_BUILD_LOG, size_ret, build_log, &size_ret); - test_error( error, "Unable to query build log" ); - log_error("ERROR: CL_PROGRAM_BUILD_LOG:\n%s\n", build_log); - free(build_log); - - } - - } - - // Was the number of devices given - if (num_devices == 0) { - - // If zero devices were specified then free the device list - free(device_list); - - } - - } - - return CL_SUCCESS; -} diff --git a/test_conformance/compatibility/test_common/harness/imageHelpers.cpp b/test_conformance/compatibility/test_common/harness/imageHelpers.cpp deleted file mode 100644 index 60f744b4..00000000 --- a/test_conformance/compatibility/test_common/harness/imageHelpers.cpp +++ /dev/null @@ -1,249 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "imageHelpers.h" - -size_t get_format_type_size( const cl_image_format *format ) -{ - return get_channel_data_type_size( format->image_channel_data_type ); -} - -size_t get_channel_data_type_size( cl_channel_type channelType ) -{ - switch( channelType ) - { - case CL_SNORM_INT8: - case CL_UNORM_INT8: - case CL_SIGNED_INT8: - case CL_UNSIGNED_INT8: - return 1; - - case CL_SNORM_INT16: - case CL_UNORM_INT16: - case CL_SIGNED_INT16: - case CL_UNSIGNED_INT16: - case CL_HALF_FLOAT: -#ifdef CL_SFIXED14_APPLE - case CL_SFIXED14_APPLE: -#endif - return sizeof( cl_short ); - - case CL_SIGNED_INT32: - case CL_UNSIGNED_INT32: - return sizeof( cl_int ); - - case CL_UNORM_SHORT_565: - case CL_UNORM_SHORT_555: -#ifdef OBSOLETE_FORAMT - case CL_UNORM_SHORT_565_REV: - case CL_UNORM_SHORT_555_REV: -#endif - return 2; - -#ifdef OBSOLETE_FORAMT - case CL_UNORM_INT_8888: - case CL_UNORM_INT_8888_REV: - return 4; -#endif - - case CL_UNORM_INT_101010: -#ifdef OBSOLETE_FORAMT - case CL_UNORM_INT_101010_REV: -#endif - return 4; - - case CL_FLOAT: - return sizeof( cl_float ); - - default: - return 0; - } -} - -size_t get_format_channel_count( const cl_image_format *format ) -{ - return get_channel_order_channel_count( format->image_channel_order ); -} - -size_t get_channel_order_channel_count( cl_channel_order order ) -{ - switch( order ) - { - case CL_R: - case CL_A: - case CL_Rx: - case CL_INTENSITY: - case CL_LUMINANCE: - return 1; - - case CL_RG: - case CL_RA: - case CL_RGx: - return 2; - - case CL_RGB: - case CL_RGBx: - return 3; - - case CL_RGBA: - case CL_ARGB: - case CL_BGRA: -#ifdef CL_1RGB_APPLE - case CL_1RGB_APPLE: -#endif -#ifdef CL_BGR1_APPLE - case CL_BGR1_APPLE: -#endif - return 4; - - default: - return 0; - } -} - -int is_format_signed( const cl_image_format *format ) -{ - switch( format->image_channel_data_type ) - { - case CL_SNORM_INT8: - case CL_SIGNED_INT8: - case CL_SNORM_INT16: - case CL_SIGNED_INT16: - case CL_SIGNED_INT32: - case CL_HALF_FLOAT: - case CL_FLOAT: -#ifdef CL_SFIXED14_APPLE - case CL_SFIXED14_APPLE: -#endif - return 1; - - default: - return 0; - } -} - -size_t get_pixel_size( cl_image_format *format ) -{ - switch( format->image_channel_data_type ) - { - case CL_SNORM_INT8: - case CL_UNORM_INT8: - case CL_SIGNED_INT8: - case CL_UNSIGNED_INT8: - return get_format_channel_count( format ); - - case CL_SNORM_INT16: - case CL_UNORM_INT16: - case CL_SIGNED_INT16: - case CL_UNSIGNED_INT16: - case CL_HALF_FLOAT: -#ifdef CL_SFIXED14_APPLE - case CL_SFIXED14_APPLE: -#endif - return get_format_channel_count( format ) * sizeof( cl_ushort ); - - case CL_SIGNED_INT32: - case CL_UNSIGNED_INT32: - return get_format_channel_count( format ) * sizeof( cl_int ); - - case CL_UNORM_SHORT_565: - case CL_UNORM_SHORT_555: -#ifdef OBSOLETE_FORAMT - case CL_UNORM_SHORT_565_REV: - case CL_UNORM_SHORT_555_REV: -#endif - return 2; - -#ifdef OBSOLETE_FORAMT - case CL_UNORM_INT_8888: - case CL_UNORM_INT_8888_REV: - return 4; -#endif - - case CL_UNORM_INT_101010: -#ifdef OBSOLETE_FORAMT - case CL_UNORM_INT_101010_REV: -#endif - return 4; - - case CL_FLOAT: - return get_format_channel_count( format ) * sizeof( cl_float ); - - default: - return 0; - } -} - -int get_8_bit_image_format( cl_context context, cl_mem_object_type objType, cl_mem_flags flags, size_t channelCount, cl_image_format *outFormat ) -{ - cl_image_format formatList[ 128 ]; - unsigned int outFormatCount, i; - int error; - - - /* Make sure each image format is supported */ - if ((error = clGetSupportedImageFormats( context, flags, objType, 128, formatList, &outFormatCount ))) - return error; - - - /* Look for one that is an 8-bit format */ - for( i = 0; i < outFormatCount; i++ ) - { - if( formatList[ i ].image_channel_data_type == CL_SNORM_INT8 || - formatList[ i ].image_channel_data_type == CL_UNORM_INT8 || - formatList[ i ].image_channel_data_type == CL_SIGNED_INT8 || - formatList[ i ].image_channel_data_type == CL_UNSIGNED_INT8 ) - { - if ( !channelCount || ( channelCount && ( get_format_channel_count( &formatList[ i ] ) == channelCount ) ) ) - { - *outFormat = formatList[ i ]; - return 0; - } - } - } - - return -1; -} - -int get_32_bit_image_format( cl_context context, cl_mem_object_type objType, cl_mem_flags flags, size_t channelCount, cl_image_format *outFormat ) -{ - cl_image_format formatList[ 128 ]; - unsigned int outFormatCount, i; - int error; - - - /* Make sure each image format is supported */ - if ((error = clGetSupportedImageFormats( context, flags, objType, 128, formatList, &outFormatCount ))) - return error; - - /* Look for one that is an 8-bit format */ - for( i = 0; i < outFormatCount; i++ ) - { - if( formatList[ i ].image_channel_data_type == CL_UNORM_INT_101010 || - formatList[ i ].image_channel_data_type == CL_FLOAT || - formatList[ i ].image_channel_data_type == CL_SIGNED_INT32 || - formatList[ i ].image_channel_data_type == CL_UNSIGNED_INT32 ) - { - if ( !channelCount || ( channelCount && ( get_format_channel_count( &formatList[ i ] ) == channelCount ) ) ) - { - *outFormat = formatList[ i ]; - return 0; - } - } - } - - return -1; -} - diff --git a/test_conformance/compatibility/test_common/harness/imageHelpers.h b/test_conformance/compatibility/test_common/harness/imageHelpers.h deleted file mode 100644 index 05b7832f..00000000 --- a/test_conformance/compatibility/test_common/harness/imageHelpers.h +++ /dev/null @@ -1,37 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#ifndef _imageHelpers_h -#define _imageHelpers_h - -#include "errorHelpers.h" - - -extern size_t get_format_type_size( const cl_image_format *format ); -extern size_t get_channel_data_type_size( cl_channel_type channelType ); -extern size_t get_format_channel_count( const cl_image_format *format ); -extern size_t get_channel_order_channel_count( cl_channel_order order ); -extern int is_format_signed( const cl_image_format *format ); -extern size_t get_pixel_size( cl_image_format *format ); - -/* Helper to get any ol image format as long as it is 8-bits-per-channel */ -extern int get_8_bit_image_format( cl_context context, cl_mem_object_type objType, cl_mem_flags flags, size_t channelCount, cl_image_format *outFormat ); - -/* Helper to get any ol image format as long as it is 32-bits-per-channel */ -extern int get_32_bit_image_format( cl_context context, cl_mem_object_type objType, cl_mem_flags flags, size_t channelCount, cl_image_format *outFormat ); - - -#endif // _imageHelpers_h - diff --git a/test_conformance/compatibility/test_common/harness/kernelHelpers.c b/test_conformance/compatibility/test_common/harness/kernelHelpers.c deleted file mode 100644 index 42f41aa4..00000000 --- a/test_conformance/compatibility/test_common/harness/kernelHelpers.c +++ /dev/null @@ -1,684 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "kernelHelpers.h" -#include "errorHelpers.h" -#include "imageHelpers.h" - -#if defined(__MINGW32__) -#include "mingw_compat.h" -#endif - -int create_single_kernel_helper( cl_context context, cl_program *outProgram, cl_kernel *outKernel, unsigned int numKernelLines, const char **kernelProgram, const char *kernelName ) -{ - int error = CL_SUCCESS; - - /* Create the program object from source */ - *outProgram = clCreateProgramWithSource( context, numKernelLines, kernelProgram, NULL, &error ); - if( *outProgram == NULL || error != CL_SUCCESS) - { - print_error( error, "clCreateProgramWithSource failed" ); - return error; - } - - /* Compile the program */ - int buildProgramFailed = 0; - int printedSource = 0; - error = clBuildProgram( *outProgram, 0, NULL, NULL, NULL, NULL ); - if (error != CL_SUCCESS) - { - unsigned int i; - print_error(error, "clBuildProgram failed"); - buildProgramFailed = 1; - printedSource = 1; - log_error( "Original source is: ------------\n" ); - for( i = 0; i < numKernelLines; i++ ) - log_error( "%s", kernelProgram[ i ] ); - } - - // Verify the build status on all devices - cl_uint deviceCount = 0; - error = clGetProgramInfo( *outProgram, CL_PROGRAM_NUM_DEVICES, sizeof( deviceCount ), &deviceCount, NULL ); - if (error != CL_SUCCESS) { - print_error(error, "clGetProgramInfo CL_PROGRAM_NUM_DEVICES failed"); - return error; - } - - if (deviceCount == 0) { - log_error("No devices found for program.\n"); - return -1; - } - - cl_device_id *devices = (cl_device_id*) malloc( deviceCount * sizeof( cl_device_id ) ); - if( NULL == devices ) - return -1; - memset( devices, 0, deviceCount * sizeof( cl_device_id )); - error = clGetProgramInfo( *outProgram, CL_PROGRAM_DEVICES, sizeof( cl_device_id ) * deviceCount, devices, NULL ); - if (error != CL_SUCCESS) { - print_error(error, "clGetProgramInfo CL_PROGRAM_DEVICES failed"); - free( devices ); - return error; - } - - cl_uint z; - for( z = 0; z < deviceCount; z++ ) - { - char deviceName[4096] = ""; - error = clGetDeviceInfo(devices[z], CL_DEVICE_NAME, sizeof( deviceName), deviceName, NULL); - if (error != CL_SUCCESS || deviceName[0] == '\0') { - log_error("Device \"%d\" failed to return a name\n", z); - print_error(error, "clGetDeviceInfo CL_DEVICE_NAME failed"); - } - - cl_build_status buildStatus; - error = clGetProgramBuildInfo(*outProgram, devices[z], CL_PROGRAM_BUILD_STATUS, sizeof(buildStatus), &buildStatus, NULL); - if (error != CL_SUCCESS) { - print_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_STATUS failed"); - free( devices ); - return error; - } - - if (buildStatus != CL_BUILD_SUCCESS || buildProgramFailed) { - char log[10240] = ""; - if (buildStatus == CL_BUILD_SUCCESS && buildProgramFailed) log_error("clBuildProgram returned an error, but buildStatus is marked as CL_BUILD_SUCCESS.\n"); - - char statusString[64] = ""; - if (buildStatus == (cl_build_status)CL_BUILD_SUCCESS) - sprintf(statusString, "CL_BUILD_SUCCESS"); - else if (buildStatus == (cl_build_status)CL_BUILD_NONE) - sprintf(statusString, "CL_BUILD_NONE"); - else if (buildStatus == (cl_build_status)CL_BUILD_ERROR) - sprintf(statusString, "CL_BUILD_ERROR"); - else if (buildStatus == (cl_build_status)CL_BUILD_IN_PROGRESS) - sprintf(statusString, "CL_BUILD_IN_PROGRESS"); - else - sprintf(statusString, "UNKNOWN (%d)", buildStatus); - - if (buildStatus != CL_BUILD_SUCCESS) log_error("Build not successful for device \"%s\", status: %s\n", deviceName, statusString); - error = clGetProgramBuildInfo( *outProgram, devices[z], CL_PROGRAM_BUILD_LOG, sizeof(log), log, NULL ); - if (error != CL_SUCCESS || log[0]=='\0'){ - log_error("Device %d (%s) failed to return a build log\n", z, deviceName); - if (error) { - print_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_LOG failed"); - free( devices ); - return error; - } else { - log_error("clGetProgramBuildInfo returned an empty log.\n"); - free( devices ); - return -1; - } - } - // In this case we've already printed out the code above. - if (!printedSource) - { - unsigned int i; - log_error( "Original source is: ------------\n" ); - for( i = 0; i < numKernelLines; i++ ) - log_error( "%s", kernelProgram[ i ] ); - printedSource = 1; - } - log_error( "Build log for device \"%s\" is: ------------\n", deviceName ); - log_error( "%s\n", log ); - log_error( "\n----------\n" ); - free( devices ); - return -1; - } - } - - /* And create a kernel from it */ - *outKernel = clCreateKernel( *outProgram, kernelName, &error ); - if( *outKernel == NULL || error != CL_SUCCESS) - { - print_error( error, "Unable to create kernel" ); - free( devices ); - return error; - } - - free( devices ); - return 0; -} - -int get_device_version( cl_device_id id, size_t* major, size_t* minor) -{ - cl_char buffer[ 4098 ]; - size_t length; - - // Device version should fit the regex "OpenCL [0-9]+\.[0-9]+ *.*" - cl_int error = clGetDeviceInfo( id, CL_DEVICE_VERSION, sizeof( buffer ), buffer, &length ); - test_error( error, "Unable to get device version string" ); - - char *p1 = (char *)buffer + strlen( "OpenCL " ); - char *p2; - while( *p1 == ' ' ) - p1++; - *major = strtol( p1, &p2, 10 ); - error = *p2 != '.'; - test_error(error, "ERROR: Version number must contain a decimal point!"); - *minor = strtol( ++p2, NULL, 10 ); - return error; -} - -int get_max_allowed_work_group_size( cl_context context, cl_kernel kernel, size_t *outMaxSize, size_t *outLimits ) -{ - cl_device_id *devices; - size_t size, maxCommonSize = 0; - int numDevices, i, j, error; - cl_uint numDims; - size_t outSize; - size_t sizeLimit[]={1,1,1}; - - - /* Assume fewer than 16 devices will be returned */ - error = clGetContextInfo( context, CL_CONTEXT_DEVICES, 0, NULL, &outSize ); - test_error( error, "Unable to obtain list of devices size for context" ); - devices = (cl_device_id *)malloc(outSize); - error = clGetContextInfo( context, CL_CONTEXT_DEVICES, outSize, devices, NULL ); - test_error( error, "Unable to obtain list of devices for context" ); - - numDevices = (int)( outSize / sizeof( cl_device_id ) ); - - for( i = 0; i < numDevices; i++ ) - { - error = clGetDeviceInfo( devices[i], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof( size ), &size, NULL ); - test_error( error, "Unable to obtain max work group size for device" ); - if( size < maxCommonSize || maxCommonSize == 0) - maxCommonSize = size; - - error = clGetKernelWorkGroupInfo( kernel, devices[i], CL_KERNEL_WORK_GROUP_SIZE, sizeof( size ), &size, NULL ); - test_error( error, "Unable to obtain max work group size for device and kernel combo" ); - if( size < maxCommonSize || maxCommonSize == 0) - maxCommonSize = size; - - error= clGetDeviceInfo( devices[i], CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof( numDims ), &numDims, NULL); - test_error( error, "clGetDeviceInfo failed for CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS"); - sizeLimit[0] = 1; - error= clGetDeviceInfo( devices[i], CL_DEVICE_MAX_WORK_ITEM_SIZES, numDims*sizeof(size_t), sizeLimit, NULL); - test_error( error, "clGetDeviceInfo failed for CL_DEVICE_MAX_WORK_ITEM_SIZES"); - - if (outLimits != NULL) - { - if (i == 0) { - for (j=0; j<3; j++) - outLimits[j] = sizeLimit[j]; - } else { - for (j=0; j<(int)numDims; j++) { - if (sizeLimit[j] < outLimits[j]) - outLimits[j] = sizeLimit[j]; - } - } - } - } - free(devices); - - *outMaxSize = (unsigned int)maxCommonSize; - return 0; -} - - -int get_max_common_work_group_size( cl_context context, cl_kernel kernel, - size_t globalThreadSize, size_t *outMaxSize ) -{ - size_t sizeLimit[3]; - int error = get_max_allowed_work_group_size( context, kernel, outMaxSize, sizeLimit ); - if( error != 0 ) - return error; - - /* Now find the largest factor of globalThreadSize that is <= maxCommonSize */ - /* Note for speed, we don't need to check the range of maxCommonSize, b/c once it gets to 1, - the modulo test will succeed and break the loop anyway */ - for( ; ( globalThreadSize % *outMaxSize ) != 0 || (*outMaxSize > sizeLimit[0]); (*outMaxSize)-- ) - ; - return 0; -} - -int get_max_common_2D_work_group_size( cl_context context, cl_kernel kernel, - size_t *globalThreadSizes, size_t *outMaxSizes ) -{ - size_t sizeLimit[3]; - size_t maxSize; - int error = get_max_allowed_work_group_size( context, kernel, &maxSize, sizeLimit ); - if( error != 0 ) - return error; - - /* Now find a set of factors, multiplied together less than maxSize, but each a factor of the global - sizes */ - - /* Simple case */ - if( globalThreadSizes[ 0 ] * globalThreadSizes[ 1 ] <= maxSize ) - { - if (globalThreadSizes[ 0 ] <= sizeLimit[0] && globalThreadSizes[ 1 ] <= sizeLimit[1]) { - outMaxSizes[ 0 ] = globalThreadSizes[ 0 ]; - outMaxSizes[ 1 ] = globalThreadSizes[ 1 ]; - return 0; - } - } - - size_t remainingSize, sizeForThisOne; - remainingSize = maxSize; - int i, j; - for (i=0 ; i<2; i++) { - if (globalThreadSizes[i] > remainingSize) - sizeForThisOne = remainingSize; - else - sizeForThisOne = globalThreadSizes[i]; - for (; (globalThreadSizes[i] % sizeForThisOne) != 0 || (sizeForThisOne > sizeLimit[i]); sizeForThisOne--) ; - outMaxSizes[i] = sizeForThisOne; - remainingSize = maxSize; - for (j=0; j<=i; j++) - remainingSize /=outMaxSizes[j]; - } - - return 0; -} - -int get_max_common_3D_work_group_size( cl_context context, cl_kernel kernel, - size_t *globalThreadSizes, size_t *outMaxSizes ) -{ - size_t sizeLimit[3]; - size_t maxSize; - int error = get_max_allowed_work_group_size( context, kernel, &maxSize, sizeLimit ); - if( error != 0 ) - return error; - /* Now find a set of factors, multiplied together less than maxSize, but each a factor of the global - sizes */ - - /* Simple case */ - if( globalThreadSizes[ 0 ] * globalThreadSizes[ 1 ] * globalThreadSizes[ 2 ] <= maxSize ) - { - if (globalThreadSizes[ 0 ] <= sizeLimit[0] && globalThreadSizes[ 1 ] <= sizeLimit[1] && globalThreadSizes[ 2 ] <= sizeLimit[2]) { - outMaxSizes[ 0 ] = globalThreadSizes[ 0 ]; - outMaxSizes[ 1 ] = globalThreadSizes[ 1 ]; - outMaxSizes[ 2 ] = globalThreadSizes[ 2 ]; - return 0; - } - } - - size_t remainingSize, sizeForThisOne; - remainingSize = maxSize; - int i, j; - for (i=0 ; i<3; i++) { - if (globalThreadSizes[i] > remainingSize) - sizeForThisOne = remainingSize; - else - sizeForThisOne = globalThreadSizes[i]; - for (; (globalThreadSizes[i] % sizeForThisOne) != 0 || (sizeForThisOne > sizeLimit[i]); sizeForThisOne--) ; - outMaxSizes[i] = sizeForThisOne; - remainingSize = maxSize; - for (j=0; j<=i; j++) - remainingSize /=outMaxSizes[j]; - } - - return 0; -} - -/* Helper to determine if an extension is supported by a device */ -int is_extension_available( cl_device_id device, const char *extensionName ) -{ - char *extString; - size_t size = 0; - int err; - int result = 0; - - if(( err = clGetDeviceInfo(device, CL_DEVICE_EXTENSIONS, 0, NULL, &size) )) - { - log_error( "Error: failed to determine size of device extensions string at %s:%d (err = %d)\n", __FILE__, __LINE__, err ); - return 0; - } - - if( 0 == size ) - return 0; - - extString = (char*) malloc( size ); - if( NULL == extString ) - { - log_error( "Error: unable to allocate %ld byte buffer for extension string at %s:%d (err = %d)\n", size, __FILE__, __LINE__, err ); - return 0; - } - - if(( err = clGetDeviceInfo(device, CL_DEVICE_EXTENSIONS, size, extString, NULL) )) - { - log_error( "Error: failed to obtain device extensions string at %s:%d (err = %d)\n", __FILE__, __LINE__, err ); - free( extString ); - return 0; - } - - if( strstr( extString, extensionName ) ) - result = 1; - - free( extString ); - return result; -} - -/* Helper to determine if a device supports an image format */ -int is_image_format_supported( cl_context context, cl_mem_flags flags, cl_mem_object_type image_type, const cl_image_format *fmt ) -{ - cl_image_format *list; - cl_uint count = 0; - cl_int err = clGetSupportedImageFormats( context, flags, image_type, 128, NULL, &count ); - if( count == 0 ) - return 0; - - list = (cl_image_format*) malloc( count * sizeof( cl_image_format ) ); - if( NULL == list ) - { - log_error( "Error: unable to allocate %ld byte buffer for image format list at %s:%d (err = %d)\n", count * sizeof( cl_image_format ), __FILE__, __LINE__, err ); - return 0; - } - - cl_int error = clGetSupportedImageFormats( context, flags, image_type, count, list, NULL ); - if( error ) - { - log_error( "Error: failed to obtain supported image type list at %s:%d (err = %d)\n", __FILE__, __LINE__, err ); - free( list ); - return 0; - } - - // iterate looking for a match. - cl_uint i; - for( i = 0; i < count; i++ ) - { - if( fmt->image_channel_data_type == list[ i ].image_channel_data_type && - fmt->image_channel_order == list[ i ].image_channel_order ) - break; - } - - free( list ); - return ( i < count ) ? true : false; -} - -size_t get_pixel_bytes( const cl_image_format *fmt ); -size_t get_pixel_bytes( const cl_image_format *fmt ) -{ - size_t chanCount; - switch( fmt->image_channel_order ) - { - case CL_R: - case CL_A: - case CL_Rx: - case CL_INTENSITY: - case CL_LUMINANCE: - chanCount = 1; - break; - case CL_RG: - case CL_RA: - case CL_RGx: - chanCount = 2; - break; - case CL_RGB: - case CL_RGBx: - chanCount = 3; - break; - case CL_RGBA: - case CL_ARGB: - case CL_BGRA: -#ifdef CL_1RGB_APPLE - case CL_1RGB_APPLE: -#endif -#ifdef CL_BGR1_APPLE - case CL_BGR1_APPLE: -#endif - chanCount = 4; - break; - default: - log_error("Unknown channel order at %s:%d!\n", __FILE__, __LINE__ ); - abort(); - break; - } - - switch( fmt->image_channel_data_type ) - { - case CL_UNORM_SHORT_565: - case CL_UNORM_SHORT_555: - return 2; - - case CL_UNORM_INT_101010: - return 4; - - case CL_SNORM_INT8: - case CL_UNORM_INT8: - case CL_SIGNED_INT8: - case CL_UNSIGNED_INT8: - return chanCount; - - case CL_SNORM_INT16: - case CL_UNORM_INT16: - case CL_HALF_FLOAT: - case CL_SIGNED_INT16: - case CL_UNSIGNED_INT16: -#ifdef CL_SFIXED14_APPLE - case CL_SFIXED14_APPLE: -#endif - return chanCount * 2; - - case CL_SIGNED_INT32: - case CL_UNSIGNED_INT32: - case CL_FLOAT: - return chanCount * 4; - - default: - log_error("Unknown channel data type at %s:%d!\n", __FILE__, __LINE__ ); - abort(); - } - - return 0; -} - -test_status verifyImageSupport( cl_device_id device ) -{ - if( checkForImageSupport( device ) ) - { - log_error( "ERROR: Device does not supported images as required by this test!\n" ); - return TEST_FAIL; - } - return TEST_PASS; -} - -int checkForImageSupport( cl_device_id device ) -{ - cl_uint i; - int error; - - - /* Check the device props to see if images are supported at all first */ - error = clGetDeviceInfo( device, CL_DEVICE_IMAGE_SUPPORT, sizeof( i ), &i, NULL ); - test_error( error, "Unable to query device for image support" ); - if( i == 0 ) - { - return CL_IMAGE_FORMAT_NOT_SUPPORTED; - } - - /* So our support is good */ - return 0; -} - -int checkFor3DImageSupport( cl_device_id device ) -{ - cl_uint i; - int error; - - /* Check the device props to see if images are supported at all first */ - error = clGetDeviceInfo( device, CL_DEVICE_IMAGE_SUPPORT, sizeof( i ), &i, NULL ); - test_error( error, "Unable to query device for image support" ); - if( i == 0 ) - { - return CL_IMAGE_FORMAT_NOT_SUPPORTED; - } - - char profile[128]; - error = clGetDeviceInfo( device, CL_DEVICE_PROFILE, sizeof(profile ), profile, NULL ); - test_error( error, "Unable to query device for CL_DEVICE_PROFILE" ); - if( 0 == strcmp( profile, "EMBEDDED_PROFILE" ) ) - { - size_t width = -1L; - size_t height = -1L; - size_t depth = -1L; - error = clGetDeviceInfo( device, CL_DEVICE_IMAGE3D_MAX_WIDTH, sizeof(width), &width, NULL ); - test_error( error, "Unable to get CL_DEVICE_IMAGE3D_MAX_WIDTH" ); - error = clGetDeviceInfo( device, CL_DEVICE_IMAGE3D_MAX_HEIGHT, sizeof(height), &height, NULL ); - test_error( error, "Unable to get CL_DEVICE_IMAGE3D_MAX_HEIGHT" ); - error = clGetDeviceInfo( device, CL_DEVICE_IMAGE3D_MAX_DEPTH, sizeof(depth), &depth, NULL ); - test_error( error, "Unable to get CL_DEVICE_IMAGE3D_MAX_DEPTH" ); - - if( 0 == (height | width | depth )) - return CL_IMAGE_FORMAT_NOT_SUPPORTED; - } - - /* So our support is good */ - return 0; -} - -void * align_malloc(size_t size, size_t alignment) -{ -#if defined(_WIN32) && defined(_MSC_VER) - return _aligned_malloc(size, alignment); -#elif defined(__linux__) || defined (linux) || defined(__APPLE__) - void * ptr = NULL; - if (0 == posix_memalign(&ptr, alignment, size)) - return ptr; - return NULL; -#elif defined(__MINGW32__) - return __mingw_aligned_malloc(size, alignment); -#else - #error "Please add support OS for aligned malloc" -#endif -} - -void align_free(void * ptr) -{ -#if defined(_WIN32) && defined(_MSC_VER) - _aligned_free(ptr); -#elif defined(__linux__) || defined (linux) || defined(__APPLE__) - return free(ptr); -#elif defined(__MINGW32__) - return __mingw_aligned_free(ptr); -#else - #error "Please add support OS for aligned free" -#endif -} - -size_t get_min_alignment(cl_context context) -{ - static cl_uint align_size = 0; - - if( 0 == align_size ) - { - cl_device_id * devices; - size_t devices_size = 0; - cl_uint result = 0; - cl_int error; - int i; - - error = clGetContextInfo (context, - CL_CONTEXT_DEVICES, - 0, - NULL, - &devices_size); - test_error_ret(error, "clGetContextInfo failed", 0); - - devices = (cl_device_id*)malloc(devices_size); - if (devices == NULL) { - print_error( error, "malloc failed" ); - return 0; - } - - error = clGetContextInfo (context, - CL_CONTEXT_DEVICES, - devices_size, - (void*)devices, - NULL); - test_error_ret(error, "clGetContextInfo failed", 0); - - for (i = 0; i < (int)(devices_size/sizeof(cl_device_id)); i++) - { - cl_uint alignment = 0; - - error = clGetDeviceInfo (devices[i], - CL_DEVICE_MEM_BASE_ADDR_ALIGN, - sizeof(cl_uint), - (void*)&alignment, - NULL); - - if (error == CL_SUCCESS) - { - alignment >>= 3; // convert bits to bytes - result = (alignment > result) ? alignment : result; - } - else - print_error( error, "clGetDeviceInfo failed" ); - } - - align_size = result; - free(devices); - } - - return align_size; -} - -cl_device_fp_config get_default_rounding_mode( cl_device_id device ) -{ - char profileStr[128] = ""; - cl_device_fp_config single = 0; - int error = clGetDeviceInfo( device, CL_DEVICE_SINGLE_FP_CONFIG, sizeof( single ), &single, NULL ); - if( error ) - test_error_ret( error, "Unable to get device CL_DEVICE_SINGLE_FP_CONFIG", 0 ); - - if( single & CL_FP_ROUND_TO_NEAREST ) - return CL_FP_ROUND_TO_NEAREST; - - if( 0 == (single & CL_FP_ROUND_TO_ZERO) ) - test_error_ret( -1, "FAILURE: device must support either CL_DEVICE_SINGLE_FP_CONFIG or CL_FP_ROUND_TO_NEAREST", 0 ); - - // Make sure we are an embedded device before allowing a pass - if( (error = clGetDeviceInfo( device, CL_DEVICE_PROFILE, sizeof( profileStr ), &profileStr, NULL ) )) - test_error_ret( error, "FAILURE: Unable to get CL_DEVICE_PROFILE", 0 ); - - if( strcmp( profileStr, "EMBEDDED_PROFILE" ) ) - test_error_ret( error, "FAILURE: non-EMBEDDED_PROFILE devices must support CL_FP_ROUND_TO_NEAREST", 0 ); - - return CL_FP_ROUND_TO_ZERO; -} - -int checkDeviceForQueueSupport( cl_device_id device, cl_command_queue_properties prop ) -{ - cl_command_queue_properties realProps; - cl_int error = clGetDeviceInfo( device, CL_DEVICE_QUEUE_PROPERTIES, sizeof( realProps ), &realProps, NULL ); - test_error_ret( error, "FAILURE: Unable to get device queue properties", 0 ); - - return ( realProps & prop ) ? 1 : 0; -} - -int printDeviceHeader( cl_device_id device ) -{ - char deviceName[ 512 ], deviceVendor[ 512 ], deviceVersion[ 512 ], cLangVersion[ 512 ]; - int error; - - error = clGetDeviceInfo( device, CL_DEVICE_NAME, sizeof( deviceName ), deviceName, NULL ); - test_error( error, "Unable to get CL_DEVICE_NAME for device" ); - - error = clGetDeviceInfo( device, CL_DEVICE_VENDOR, sizeof( deviceVendor ), deviceVendor, NULL ); - test_error( error, "Unable to get CL_DEVICE_VENDOR for device" ); - - error = clGetDeviceInfo( device, CL_DEVICE_VERSION, sizeof( deviceVersion ), deviceVersion, NULL ); - test_error( error, "Unable to get CL_DEVICE_VERSION for device" ); - - error = clGetDeviceInfo( device, CL_DEVICE_OPENCL_C_VERSION, sizeof( cLangVersion ), cLangVersion, NULL ); - test_error( error, "Unable to get CL_DEVICE_OPENCL_C_VERSION for device" ); - - log_info("Compute Device Name = %s, Compute Device Vendor = %s, Compute Device Version = %s%s%s\n", - deviceName, deviceVendor, deviceVersion, ( error == CL_SUCCESS ) ? ", CL C Version = " : "", - ( error == CL_SUCCESS ) ? cLangVersion : "" ); - - return CL_SUCCESS; -} diff --git a/test_conformance/compatibility/test_common/harness/kernelHelpers.h b/test_conformance/compatibility/test_common/harness/kernelHelpers.h deleted file mode 100644 index 09515e28..00000000 --- a/test_conformance/compatibility/test_common/harness/kernelHelpers.h +++ /dev/null @@ -1,129 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#ifndef _kernelHelpers_h -#define _kernelHelpers_h - -#include "compat.h" -#include "testHarness.h" - -#include -#include - -#if defined (__MINGW32__) -#include -#endif - -#include - -#ifdef __APPLE__ - #include -#else - #include -#endif - -#ifdef __cplusplus -extern "C" { -#endif // __cplusplus - -/* - * The below code is intended to be used at the top of kernels that appear inline in files to set line and file info for the kernel: - * - * const char *source = { - * INIT_OPENCL_DEBUG_INFO - * "__kernel void foo( int x )\n" - * "{\n" - * " ...\n" - * "}\n" - * }; - */ -#define INIT_OPENCL_DEBUG_INFO SET_OPENCL_LINE_INFO( __LINE__, __FILE__ ) -#define SET_OPENCL_LINE_INFO(_line, _file) "#line " STRINGIFY(_line) " " STRINGIFY(_file) "\n" -#ifndef STRINGIFY_VALUE - #define STRINGIFY_VALUE(_x) STRINGIFY(_x) -#endif -#ifndef STRINGIFY - #define STRINGIFY(_x) #_x -#endif - -/* Helper that creates a single program and kernel from a single-kernel program source */ -extern int create_single_kernel_helper( cl_context context, cl_program *outProgram, cl_kernel *outKernel, unsigned int numKernelLines, const char **kernelProgram, const char *kernelName ); - -/* Helper to obtain the biggest fit work group size for all the devices in a given group and for the given global thread size */ -extern int get_max_common_work_group_size( cl_context context, cl_kernel kernel, size_t globalThreadSize, size_t *outSize ); - -/* Helper to obtain the biggest fit work group size for all the devices in a given group and for the given global thread size */ -extern int get_max_common_2D_work_group_size( cl_context context, cl_kernel kernel, size_t *globalThreadSize, size_t *outSizes ); - -/* Helper to obtain the biggest fit work group size for all the devices in a given group and for the given global thread size */ -extern int get_max_common_3D_work_group_size( cl_context context, cl_kernel kernel, size_t *globalThreadSize, size_t *outSizes ); - -/* Helper to get major/minor number for a device */ -extern int get_device_version( cl_device_id id, size_t* major, size_t* minor); - -/* Helper to obtain the biggest allowed work group size for all the devices in a given group */ -extern int get_max_allowed_work_group_size( cl_context context, cl_kernel kernel, size_t *outSize, size_t *outLimits ); - -/* Helper to determine if an extension is supported by a device */ -extern int is_extension_available( cl_device_id device, const char *extensionName ); - -/* Helper to determine if a device supports an image format */ -extern int is_image_format_supported( cl_context context, cl_mem_flags flags, cl_mem_object_type image_type, const cl_image_format *fmt ); - -/* Helper to get pixel size for a pixel format */ -size_t get_pixel_bytes( const cl_image_format *fmt ); - -/* Verify the given device supports images. */ -extern test_status verifyImageSupport( cl_device_id device ); - -/* Checks that the given device supports images. Same as verify, but doesn't print an error */ -extern int checkForImageSupport( cl_device_id device ); -extern int checkFor3DImageSupport( cl_device_id device ); - -/* Checks that a given queue property is supported on the specified device. Returns 1 if supported, 0 if not or an error. */ -extern int checkDeviceForQueueSupport( cl_device_id device, cl_command_queue_properties prop ); - -/* Helper for aligned memory allocation */ -void * align_malloc(size_t size, size_t alignment); -void align_free(void *); - -/* Helper to obtain the min alignment for a given context, i.e the max of all min alignments for devices attached to the context*/ -size_t get_min_alignment(cl_context context); - -/* Helper to obtain the default rounding mode for single precision computation. (Double is always CL_FP_ROUND_TO_NEAREST.) Returns 0 on error. */ -cl_device_fp_config get_default_rounding_mode( cl_device_id device ); - -#define PASSIVE_REQUIRE_IMAGE_SUPPORT( device ) \ - if( checkForImageSupport( device ) ) \ - { \ - log_info( "\n\tNote: device does not support images. Skipping test...\n" ); \ - return 0; \ - } - -#define PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( device ) \ - if( checkFor3DImageSupport( device ) ) \ - { \ - log_info( "\n\tNote: device does not support 3D images. Skipping test...\n" ); \ - return 0; \ - } - -/* Prints out the standard device header for all tests given the device to print for */ -extern int printDeviceHeader( cl_device_id device ); - -#ifdef __cplusplus -} -#endif // __cplusplus - -#endif // _kernelHelpers_h diff --git a/test_conformance/compatibility/test_common/harness/testHarness.c b/test_conformance/compatibility/test_common/harness/testHarness.c deleted file mode 100644 index f4cfb00b..00000000 --- a/test_conformance/compatibility/test_common/harness/testHarness.c +++ /dev/null @@ -1,932 +0,0 @@ -// -// Copyright (c) 2017-2019 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "testHarness.h" -#include "compat.h" -#include -#include - -#if !defined(_WIN32) -#include -#endif - -#include -#include -#include -#include -#include "threadTesting.h" -#include "errorHelpers.h" -#include "kernelHelpers.h" -#include "fpcontrol.h" - -#if !defined(_WIN32) -#include -#endif - -#include - -#if !defined (__APPLE__) -#include -#endif - -int gTestsPassed = 0; -int gTestsFailed = 0; -int gFailCount; -int gTestCount; -cl_uint gRandomSeed = 0; -cl_uint gReSeed = 0; - -int gFlushDenormsToZero = 0; -int gInfNanSupport = 1; -int gIsEmbedded = 0; -int gIsOpenCL_C_1_0_Device = 0; -int gIsOpenCL_1_0_Device = 0; -int gHasLong = 1; - -#define DEFAULT_NUM_ELEMENTS 0x4000 - -int runTestHarness( int argc, const char *argv[], int testNum, test_definition testList[], - int imageSupportRequired, int forceNoContextCreation, cl_command_queue_properties queueProps ) -{ - return runTestHarnessWithCheck( argc, argv, testNum, testList, imageSupportRequired, forceNoContextCreation, queueProps, - ( imageSupportRequired ) ? verifyImageSupport : NULL ); -} - -int runTestHarnessWithCheck( int argc, const char *argv[], int testNum, test_definition testList[], - int imageSupportRequired, int forceNoContextCreation, cl_command_queue_properties queueProps, - DeviceCheckFn deviceCheckFn ) -{ - test_start(); - log_info("*** Compatibility with Previous Versions test ***\n"); - - cl_device_type device_type = CL_DEVICE_TYPE_DEFAULT; - cl_uint num_platforms = 0; - cl_platform_id *platforms; - cl_device_id device; - int num_elements = DEFAULT_NUM_ELEMENTS; - cl_uint num_devices = 0; - cl_device_id *devices = NULL; - cl_uint choosen_device_index = 0; - cl_uint choosen_platform_index = 0; - - int err, ret; - char *endPtr; - int based_on_env_var = 0; - - - /* Check for environment variable to set device type */ - char *env_mode = getenv( "CL_DEVICE_TYPE" ); - if( env_mode != NULL ) - { - based_on_env_var = 1; - if( strcmp( env_mode, "gpu" ) == 0 || strcmp( env_mode, "CL_DEVICE_TYPE_GPU" ) == 0 ) - device_type = CL_DEVICE_TYPE_GPU; - else if( strcmp( env_mode, "cpu" ) == 0 || strcmp( env_mode, "CL_DEVICE_TYPE_CPU" ) == 0 ) - device_type = CL_DEVICE_TYPE_CPU; - else if( strcmp( env_mode, "accelerator" ) == 0 || strcmp( env_mode, "CL_DEVICE_TYPE_ACCELERATOR" ) == 0 ) - device_type = CL_DEVICE_TYPE_ACCELERATOR; - else if( strcmp( env_mode, "default" ) == 0 || strcmp( env_mode, "CL_DEVICE_TYPE_DEFAULT" ) == 0 ) - device_type = CL_DEVICE_TYPE_DEFAULT; - else - { - log_error( "Unknown CL_DEVICE_TYPE env variable setting: %s.\nAborting...\n", env_mode ); - abort(); - } - } - -#if defined( __APPLE__ ) - { - // report on any unusual library search path indirection - char *libSearchPath = getenv( "DYLD_LIBRARY_PATH"); - if( libSearchPath ) - log_info( "*** DYLD_LIBRARY_PATH = \"%s\"\n", libSearchPath ); - - // report on any unusual framework search path indirection - char *frameworkSearchPath = getenv( "DYLD_FRAMEWORK_PATH"); - if( libSearchPath ) - log_info( "*** DYLD_FRAMEWORK_PATH = \"%s\"\n", frameworkSearchPath ); - } -#endif - - env_mode = getenv( "CL_DEVICE_INDEX" ); - if( env_mode != NULL ) - { - choosen_device_index = atoi(env_mode); - } - - env_mode = getenv( "CL_PLATFORM_INDEX" ); - if( env_mode != NULL ) - { - choosen_platform_index = atoi(env_mode); - } - - /* Process the command line arguments */ - - /* Special case: just list the tests */ - if( ( argc > 1 ) && (!strcmp( argv[ 1 ], "-list" ) || !strcmp( argv[ 1 ], "-h" ) || !strcmp( argv[ 1 ], "--help" ))) - { - char *fileName = getenv("CL_CONFORMANCE_RESULTS_FILENAME"); - - log_info( "Usage: %s [*] [pid] [id] []\n", argv[0] ); - log_info( "\t\tOne or more of: (wildcard character '*') (default *)\n"); - log_info( "\tpid\tIndicates platform at index should be used (default 0).\n" ); - log_info( "\tid\t\tIndicates device at index should be used (default 0).\n" ); - log_info( "\t\tcpu|gpu|accelerator| (default CL_DEVICE_TYPE_DEFAULT)\n" ); - log_info( "\n" ); - log_info( "\tNOTE: You may pass environment variable CL_CONFORMANCE_RESULTS_FILENAME (currently '%s')\n", - fileName != NULL ? fileName : "" ); - log_info( "\t to save results to JSON file.\n" ); - - log_info( "\n" ); - log_info( "Test names:\n" ); - for( int i = 0; i < testNum; i++ ) - { - log_info( "\t%s\n", testList[i].name ); - } - test_finish(); - return EXIT_SUCCESS; - } - - /* How are we supposed to seed the random # generators? */ - if( argc > 1 && strcmp( argv[ argc - 1 ], "randomize" ) == 0 ) - { - log_info(" Initializing random seed based on the clock.\n"); - gRandomSeed = (unsigned)clock(); - gReSeed = 1; - argc--; - } - else - { - log_info(" Initializing random seed to 0.\n"); - } - - /* Do we have an integer to specify the number of elements to pass to tests? */ - if( argc > 1 ) - { - ret = (int)strtol( argv[ argc - 1 ], &endPtr, 10 ); - if( endPtr != argv[ argc - 1 ] && *endPtr == 0 ) - { - /* By spec, this means the entire string was a valid integer, so we treat it as a num_elements spec */ - /* (hence why we stored the result in ret first) */ - num_elements = ret; - log_info( "Testing with num_elements of %d\n", num_elements ); - argc--; - } - } - - /* Do we have a CPU/GPU specification? */ - if( argc > 1 ) - { - if( strcmp( argv[ argc - 1 ], "gpu" ) == 0 || strcmp( argv[ argc - 1 ], "CL_DEVICE_TYPE_GPU" ) == 0 ) - { - device_type = CL_DEVICE_TYPE_GPU; - argc--; - } - else if( strcmp( argv[ argc - 1 ], "cpu" ) == 0 || strcmp( argv[ argc - 1 ], "CL_DEVICE_TYPE_CPU" ) == 0 ) - { - device_type = CL_DEVICE_TYPE_CPU; - argc--; - } - else if( strcmp( argv[ argc - 1 ], "accelerator" ) == 0 || strcmp( argv[ argc - 1 ], "CL_DEVICE_TYPE_ACCELERATOR" ) == 0 ) - { - device_type = CL_DEVICE_TYPE_ACCELERATOR; - argc--; - } - else if( strcmp( argv[ argc - 1 ], "CL_DEVICE_TYPE_DEFAULT" ) == 0 ) - { - device_type = CL_DEVICE_TYPE_DEFAULT; - argc--; - } - } - - /* Did we choose a specific device index? */ - if( argc > 1 ) - { - if( strlen( argv[ argc - 1 ] ) >= 3 && argv[ argc - 1 ][0] == 'i' && argv[ argc - 1 ][1] == 'd' ) - { - choosen_device_index = atoi( &(argv[ argc - 1 ][2]) ); - argc--; - } - } - - /* Did we choose a specific platform index? */ - if( argc > 1 ) - { - if( strlen( argv[ argc - 1 ] ) >= 3 && argv[ argc - 1 ][0] == 'p' && argv[ argc - 1 ][1] == 'i' && argv[ argc - 1 ][2] == 'd') - { - choosen_platform_index = atoi( &(argv[ argc - 1 ][3]) ); - argc--; - } - } - - switch( device_type ) - { - case CL_DEVICE_TYPE_GPU: log_info( "Requesting GPU device " ); break; - case CL_DEVICE_TYPE_CPU: log_info( "Requesting CPU device " ); break; - case CL_DEVICE_TYPE_ACCELERATOR: log_info( "Requesting Accelerator device " ); break; - case CL_DEVICE_TYPE_DEFAULT: log_info( "Requesting Default device " ); break; - default: log_error( "Requesting unknown device "); return EXIT_FAILURE; - } - log_info( based_on_env_var ? "based on environment variable " : "based on command line " ); - log_info( "for platform index %d and device index %d\n", choosen_platform_index, choosen_device_index); - -#if defined( __APPLE__ ) -#if defined( __i386__ ) || defined( __x86_64__ ) -#define kHasSSE3 0x00000008 -#define kHasSupplementalSSE3 0x00000100 -#define kHasSSE4_1 0x00000400 -#define kHasSSE4_2 0x00000800 - /* check our environment for a hint to disable SSE variants */ - { - const char *env = getenv( "CL_MAX_SSE" ); - if( env ) - { - extern int _cpu_capabilities; - int mask = 0; - if( 0 == strcasecmp( env, "SSE4.1" ) ) - mask = kHasSSE4_2; - else if( 0 == strcasecmp( env, "SSSE3" ) ) - mask = kHasSSE4_2 | kHasSSE4_1; - else if( 0 == strcasecmp( env, "SSE3" ) ) - mask = kHasSSE4_2 | kHasSSE4_1 | kHasSupplementalSSE3; - else if( 0 == strcasecmp( env, "SSE2" ) ) - mask = kHasSSE4_2 | kHasSSE4_1 | kHasSupplementalSSE3 | kHasSSE3; - else - { - log_error( "Error: Unknown CL_MAX_SSE setting: %s\n", env ); - return EXIT_FAILURE; - } - - log_info( "*** Environment: CL_MAX_SSE = %s ***\n", env ); - _cpu_capabilities &= ~mask; - } - } -#endif -#endif - - /* Get the platform */ - err = clGetPlatformIDs(0, NULL, &num_platforms); - if (err) { - print_error(err, "clGetPlatformIDs failed"); - test_finish(); - return EXIT_FAILURE; - } - - platforms = (cl_platform_id *) malloc( num_platforms * sizeof( cl_platform_id ) ); - if (!platforms || choosen_platform_index >= num_platforms) { - log_error( "platform index out of range -- choosen_platform_index (%d) >= num_platforms (%d)\n", choosen_platform_index, num_platforms ); - test_finish(); - return EXIT_FAILURE; - } - - err = clGetPlatformIDs(num_platforms, platforms, NULL); - if (err) { - print_error(err, "clGetPlatformIDs failed"); - test_finish(); - return EXIT_FAILURE; - } - - /* Get the number of requested devices */ - err = clGetDeviceIDs(platforms[choosen_platform_index], device_type, 0, NULL, &num_devices ); - if (err) { - print_error(err, "clGetDeviceIDs failed"); - test_finish(); - return EXIT_FAILURE; - } - - devices = (cl_device_id *) malloc( num_devices * sizeof( cl_device_id ) ); - if (!devices || choosen_device_index >= num_devices) { - log_error( "device index out of range -- choosen_device_index (%d) >= num_devices (%d)\n", choosen_device_index, num_devices ); - test_finish(); - return EXIT_FAILURE; - } - - /* Get the requested device */ - err = clGetDeviceIDs(platforms[choosen_platform_index], device_type, num_devices, devices, NULL ); - if (err) { - print_error(err, "clGetDeviceIDs failed"); - test_finish(); - return EXIT_FAILURE; - } - - device = devices[choosen_device_index]; - free(devices); - devices = NULL; - free(platforms); - platforms = NULL; - - if( printDeviceHeader( device ) != CL_SUCCESS ) - { - test_finish(); - return EXIT_FAILURE; - } - - cl_device_fp_config fpconfig = 0; - err = clGetDeviceInfo( device, CL_DEVICE_SINGLE_FP_CONFIG, sizeof( fpconfig ), &fpconfig, NULL ); - if (err) { - print_error(err, "clGetDeviceInfo for CL_DEVICE_SINGLE_FP_CONFIG failed"); - test_finish(); - return EXIT_FAILURE; - } - - gFlushDenormsToZero = ( 0 == (fpconfig & CL_FP_DENORM)); - log_info( "Supports single precision denormals: %s\n", gFlushDenormsToZero ? "NO" : "YES" ); - log_info( "sizeof( void*) = %d (host)\n", (int) sizeof( void* ) ); - - //detect whether profile of the device is embedded - char profile[1024] = ""; - err = clGetDeviceInfo(device, CL_DEVICE_PROFILE, sizeof(profile), profile, NULL); - if (err) - { - print_error(err, "clGetDeviceInfo for CL_DEVICE_PROFILE failed\n" ); - test_finish(); - return EXIT_FAILURE; - } - gIsEmbedded = NULL != strstr(profile, "EMBEDDED_PROFILE"); - - //detect the floating point capabilities - cl_device_fp_config floatCapabilities = 0; - err = clGetDeviceInfo(device, CL_DEVICE_SINGLE_FP_CONFIG, sizeof(floatCapabilities), &floatCapabilities, NULL); - if (err) - { - print_error(err, "clGetDeviceInfo for CL_DEVICE_SINGLE_FP_CONFIG failed\n"); - test_finish(); - return EXIT_FAILURE; - } - - // Check for problems that only embedded will have - if( gIsEmbedded ) - { - //If the device is embedded, we need to detect if the device supports Infinity and NaN - if ((floatCapabilities & CL_FP_INF_NAN) == 0) - gInfNanSupport = 0; - - // check the extensions list to see if ulong and long are supported - size_t extensionsStringSize = 0; - if( (err = clGetDeviceInfo( device, CL_DEVICE_EXTENSIONS, 0, NULL, &extensionsStringSize ) )) - { - print_error( err, "Unable to get extensions string size for embedded device" ); - test_finish(); - return EXIT_FAILURE; - } - char *extensions_string = (char*) malloc(extensionsStringSize); - if( NULL == extensions_string ) - { - print_error( CL_OUT_OF_HOST_MEMORY, "Unable to allocate storage for extensions string for embedded device" ); - test_finish(); - return EXIT_FAILURE; - } - - if( (err = clGetDeviceInfo( device, CL_DEVICE_EXTENSIONS, extensionsStringSize, extensions_string, NULL ) )) - { - print_error( err, "Unable to get extensions string for embedded device" ); - test_finish(); - return EXIT_FAILURE; - } - - if( extensions_string[extensionsStringSize-1] != '\0' ) - { - log_error( "FAILURE: extensions string for embedded device is not NUL terminated" ); - test_finish(); - return EXIT_FAILURE; - } - - if( NULL == strstr( extensions_string, "cles_khr_int64" )) - gHasLong = 0; - - free(extensions_string); - } - - if( getenv( "OPENCL_1_0_DEVICE" ) ) - { - char c_version[1024]; - gIsOpenCL_1_0_Device = 1; - memset( c_version, 0, sizeof( c_version ) ); - - if( (err = clGetDeviceInfo( device, CL_DEVICE_OPENCL_C_VERSION, sizeof(c_version), c_version, NULL )) ) - { - log_error( "FAILURE: unable to get CL_DEVICE_OPENCL_C_VERSION on 1.0 device. (%d)\n", err ); - test_finish(); - return EXIT_FAILURE; - } - - if( 0 == strncmp( c_version, "OpenCL C 1.0 ", strlen( "OpenCL C 1.0 " ) ) ) - { - gIsOpenCL_C_1_0_Device = 1; - log_info( "Device is a OpenCL C 1.0 device\n" ); - } - else - log_info( "Device is a OpenCL 1.0 device, but supports OpenCL C 1.1\n" ); - } - - cl_uint device_address_bits = 0; - if( (err = clGetDeviceInfo( device, CL_DEVICE_ADDRESS_BITS, sizeof( device_address_bits ), &device_address_bits, NULL ) )) - { - print_error( err, "Unable to obtain device address bits" ); - test_finish(); - return EXIT_FAILURE; - } - if( device_address_bits ) - log_info( "sizeof( void*) = %d (device)\n", device_address_bits/8 ); - else - { - log_error("Invalid device address bit size returned by device.\n"); - test_finish(); - return EXIT_FAILURE; - } - - - /* If we have a device checking function, run it */ - if( ( deviceCheckFn != NULL ) ) - { - test_status status = deviceCheckFn( device ); - switch (status) - { - case TEST_PASS: - break; - case TEST_FAIL: - return EXIT_FAILURE; - case TEST_SKIP: - return EXIT_SUCCESS; - } - } - - if (num_elements <= 0) - num_elements = DEFAULT_NUM_ELEMENTS; - - // On most platforms which support denorm, default is FTZ off. However, - // on some hardware where the reference is computed, default might be flush denorms to zero e.g. arm. - // This creates issues in result verification. Since spec allows the implementation to either flush or - // not flush denorms to zero, an implementation may choose not be flush i.e. return denorm result whereas - // reference result may be zero (flushed denorm). Hence we need to disable denorm flushing on host side - // where reference is being computed to make sure we get non-flushed reference result. If implementation - // returns flushed result, we correctly take care of that in verification code. -#if defined(__APPLE__) && defined(__arm__) - FPU_mode_type oldMode; - DisableFTZ( &oldMode ); -#endif - - int error = parseAndCallCommandLineTests( argc, argv, device, testNum, testList, forceNoContextCreation, queueProps, num_elements ); - - #if defined(__APPLE__) && defined(__arm__) - // Restore the old FP mode before leaving. - RestoreFPState( &oldMode ); -#endif - - return (error == 0) ? EXIT_SUCCESS : EXIT_FAILURE; -} - -static int find_matching_tests( test_definition testList[], unsigned char selectedTestList[], int testNum, - const char *argument, bool isWildcard ) -{ - int found_tests = 0; - size_t wildcard_length = strlen( argument ) - 1; /* -1 for the asterisk */ - - for( int i = 0; i < testNum; i++ ) - { - if( ( !isWildcard && strcmp( testList[i].name, argument ) == 0 ) || - ( isWildcard && strncmp( testList[i].name, argument, wildcard_length ) == 0 ) ) - { - if( selectedTestList[i] ) - { - log_error( "ERROR: Test '%s' has already been selected.\n", testList[i].name ); - return EXIT_FAILURE; - } - else if( testList[i].func == NULL ) - { - log_error( "ERROR: Test '%s' is missing implementation.\n", testList[i].name ); - return EXIT_FAILURE; - } - else - { - selectedTestList[i] = 1; - found_tests = 1; - if( !isWildcard ) - { - break; - } - } - } - } - - if( !found_tests ) - { - log_error( "ERROR: The argument '%s' did not match any test names.\n", argument ); - return EXIT_FAILURE; - } - - return EXIT_SUCCESS; -} - -static int saveResultsToJson( const char *fileName, const char *suiteName, test_definition testList[], - unsigned char selectedTestList[], test_status resultTestList[], int testNum ) -{ - FILE *file = fopen( fileName, "w" ); - if( NULL == file ) - { - log_error( "ERROR: Failed to open '%s' for writing results.\n", fileName ); - return EXIT_FAILURE; - } - - const char *save_map[] = { "success", "failure" }; - const char *result_map[] = { "pass", "fail", "skip" }; - const char *linebreak[] = { "", ",\n" }; - int add_linebreak = 0; - - fprintf( file, "{\n" ); - fprintf( file, "\t\"cmd\": \"%s\",\n", suiteName ); - fprintf( file, "\t\"results\": {\n" ); - - for( int i = 0; i < testNum; ++i ) - { - if( selectedTestList[i] ) - { - fprintf( file, "%s\t\t\"%s\": \"%s\"", linebreak[add_linebreak], testList[i].name, result_map[(int)resultTestList[i]] ); - add_linebreak = 1; - } - } - fprintf( file, "\n"); - - fprintf( file, "\t}\n" ); - fprintf( file, "}\n" ); - - int ret = fclose( file ) ? 1 : 0; - - log_info( "Saving results to %s: %s!\n", fileName, save_map[ret] ); - - return ret; -} - -static void print_results( int failed, int count, const char* name ) -{ - if( count < failed ) - { - count = failed; - } - - if( failed == 0 ) - { - if( count > 1 ) - { - log_info( "PASSED %d of %d %ss.\n", count, count, name ); - } - else - { - log_info( "PASSED %s.\n", name ); - } - } - else if( failed > 0 ) - { - if( count > 1 ) - { - log_error( "FAILED %d of %d %ss.\n", failed, count, name ); - } - else - { - log_error( "FAILED %s.\n", name ); - } - } -} - -int parseAndCallCommandLineTests( int argc, const char *argv[], cl_device_id device, int testNum, - test_definition testList[], int forceNoContextCreation, - cl_command_queue_properties queueProps, int num_elements ) -{ - int ret = EXIT_SUCCESS; - - unsigned char *selectedTestList = ( unsigned char* ) calloc( testNum, 1 ); - test_status *resultTestList = NULL; - - if( argc == 1 ) - { - /* No actual arguments, all tests will be run. */ - memset( selectedTestList, 1, testNum ); - } - else - { - for( int i = 1; i < argc; i++ ) - { - if( strchr( argv[i], '*' ) != NULL ) - { - ret = find_matching_tests( testList, selectedTestList, testNum, argv[i], true ); - } - else - { - if( strcmp( argv[i], "all" ) == 0 ) - { - memset( selectedTestList, 1, testNum ); - break; - } - else - { - ret = find_matching_tests( testList, selectedTestList, testNum, argv[i], false ); - } - } - - if( ret == EXIT_FAILURE ) - { - break; - } - } - } - - if( ret == EXIT_SUCCESS ) - { - resultTestList = ( test_status* ) calloc( testNum, sizeof(*resultTestList) ); - - callTestFunctions( testList, selectedTestList, resultTestList, testNum, device, - forceNoContextCreation, num_elements, queueProps ); - - print_results( gFailCount, gTestCount, "sub-test" ); - print_results( gTestsFailed, gTestsFailed + gTestsPassed, "test" ); - - char *filename = getenv( "CL_CONFORMANCE_RESULTS_FILENAME" ); - if( filename != NULL ) - { - ret = saveResultsToJson( filename, argv[0], testList, selectedTestList, resultTestList, testNum ); - } - } - - test_finish(); - - free( selectedTestList ); - free( resultTestList ); - - return ret; -} - -void callTestFunctions( test_definition testList[], unsigned char selectedTestList[], test_status resultTestList[], - int testNum, cl_device_id deviceToUse, int forceNoContextCreation, int numElementsToUse, - cl_command_queue_properties queueProps ) -{ - for( int i = 0; i < testNum; ++i ) - { - if( selectedTestList[i] ) - { - resultTestList[i] = callSingleTestFunction( testList[i], deviceToUse, forceNoContextCreation, - numElementsToUse, queueProps ); - } - } -} - -void CL_CALLBACK notify_callback(const char *errinfo, const void *private_info, size_t cb, void *user_data) -{ - log_info( "%s\n", errinfo ); -} - -// Actual function execution -test_status callSingleTestFunction( test_definition test, cl_device_id deviceToUse, int forceNoContextCreation, - int numElementsToUse, cl_command_queue_properties queueProps ) -{ - test_status status; - cl_int error; - cl_context context = NULL; - cl_command_queue queue = NULL; - - /* Create a context to work with, unless we're told not to */ - if( !forceNoContextCreation ) - { - context = clCreateContext(NULL, 1, &deviceToUse, notify_callback, NULL, &error ); - if (!context) - { - print_error( error, "Unable to create testing context" ); - return TEST_FAIL; - } - - queue = clCreateCommandQueue( context, deviceToUse, queueProps, &error ); - if( queue == NULL ) - { - print_error( error, "Unable to create testing command queue" ); - return TEST_FAIL; - } - } - - /* Run the test and print the result */ - log_info( "%s...\n", test.name ); - fflush( stdout ); - - const Version device_version = get_device_cl_version(deviceToUse); - if (test.min_version > device_version) - { - log_info("%s skipped (requires at least version %s, but the device reports version %s)\n", - test.name, test.min_version.to_string().c_str(), device_version.to_string().c_str()); - return TEST_SKIP; - } - - if( test.func == NULL ) - { - // Skip unimplemented test, can happen when all of the tests are selected - log_info("%s test currently not implemented\n", test.name); - status = TEST_SKIP; - } - else - { - int ret = test.func(deviceToUse, context, queue, numElementsToUse); //test_threaded_function( ptr_basefn_list[i], group, context, num_elements); - if( ret == TEST_NOT_IMPLEMENTED ) - { - /* Tests can also let us know they're not implemented yet */ - log_info("%s test currently not implemented\n", test.name); - status = TEST_SKIP; - } - else - { - /* Print result */ - if( ret == 0 ) { - log_info( "%s passed\n", test.name ); - gTestsPassed++; - status = TEST_PASS; - } - else - { - log_error( "%s FAILED\n", test.name ); - gTestsFailed++; - status = TEST_FAIL; - } - } - } - - /* Release the context */ - if( !forceNoContextCreation ) - { - int error = clFinish(queue); - if (error) { - log_error("clFinish failed: %d", error); - status = TEST_FAIL; - } - clReleaseCommandQueue( queue ); - clReleaseContext( context ); - } - - return status; -} - -void checkDeviceTypeOverride( cl_device_type *inOutType ) -{ - /* Check if we are forced to CPU mode */ - char *force_cpu = getenv( "CL_DEVICE_TYPE" ); - if( force_cpu != NULL ) - { - if( strcmp( force_cpu, "gpu" ) == 0 || strcmp( force_cpu, "CL_DEVICE_TYPE_GPU" ) == 0 ) - *inOutType = CL_DEVICE_TYPE_GPU; - else if( strcmp( force_cpu, "cpu" ) == 0 || strcmp( force_cpu, "CL_DEVICE_TYPE_CPU" ) == 0 ) - *inOutType = CL_DEVICE_TYPE_CPU; - else if( strcmp( force_cpu, "accelerator" ) == 0 || strcmp( force_cpu, "CL_DEVICE_TYPE_ACCELERATOR" ) == 0 ) - *inOutType = CL_DEVICE_TYPE_ACCELERATOR; - else if( strcmp( force_cpu, "CL_DEVICE_TYPE_DEFAULT" ) == 0 ) - *inOutType = CL_DEVICE_TYPE_DEFAULT; - } - - switch( *inOutType ) - { - case CL_DEVICE_TYPE_GPU: log_info( "Requesting GPU device " ); break; - case CL_DEVICE_TYPE_CPU: log_info( "Requesting CPU device " ); break; - case CL_DEVICE_TYPE_ACCELERATOR: log_info( "Requesting Accelerator device " ); break; - case CL_DEVICE_TYPE_DEFAULT: log_info( "Requesting Default device " ); break; - default: break; - } - log_info( force_cpu != NULL ? "based on environment variable\n" : "based on command line\n" ); - -#if defined( __APPLE__ ) - { - // report on any unusual library search path indirection - char *libSearchPath = getenv( "DYLD_LIBRARY_PATH"); - if( libSearchPath ) - log_info( "*** DYLD_LIBRARY_PATH = \"%s\"\n", libSearchPath ); - - // report on any unusual framework search path indirection - char *frameworkSearchPath = getenv( "DYLD_FRAMEWORK_PATH"); - if( libSearchPath ) - log_info( "*** DYLD_FRAMEWORK_PATH = \"%s\"\n", frameworkSearchPath ); - } -#endif - -} - -#if ! defined( __APPLE__ ) -void memset_pattern4(void *dest, const void *src_pattern, size_t bytes ) -{ - uint32_t pat = ((uint32_t*) src_pattern)[0]; - size_t count = bytes / 4; - size_t i; - uint32_t *d = (uint32_t*)dest; - - for( i = 0; i < count; i++ ) - d[i] = pat; - - d += i; - - bytes &= 3; - if( bytes ) - memcpy( d, src_pattern, bytes ); -} -#endif - -extern cl_device_type GetDeviceType( cl_device_id d ) -{ - cl_device_type result = -1; - cl_int err = clGetDeviceInfo( d, CL_DEVICE_TYPE, sizeof( result ), &result, NULL ); - if( CL_SUCCESS != err ) - log_error( "ERROR: Unable to get device type for device %p\n", d ); - return result; -} - - -cl_device_id GetOpposingDevice( cl_device_id device ) -{ - cl_int error; - cl_device_id *otherDevices; - cl_uint actualCount; - cl_platform_id plat; - - // Get the platform of the device to use for getting a list of devices - error = clGetDeviceInfo( device, CL_DEVICE_PLATFORM, sizeof( plat ), &plat, NULL ); - if( error != CL_SUCCESS ) - { - print_error( error, "Unable to get device's platform" ); - return NULL; - } - - // Get a list of all devices - error = clGetDeviceIDs( plat, CL_DEVICE_TYPE_ALL, 0, NULL, &actualCount ); - if( error != CL_SUCCESS ) - { - print_error( error, "Unable to get list of devices size" ); - return NULL; - } - otherDevices = (cl_device_id *)malloc(actualCount*sizeof(cl_device_id)); - error = clGetDeviceIDs( plat, CL_DEVICE_TYPE_ALL, actualCount, otherDevices, NULL ); - if( error != CL_SUCCESS ) - { - print_error( error, "Unable to get list of devices" ); - free(otherDevices); - return NULL; - } - - if( actualCount == 1 ) - { - free(otherDevices); - return device; // NULL means error, returning self means we couldn't find another one - } - - // Loop and just find one that isn't the one we were given - cl_uint i; - for( i = 0; i < actualCount; i++ ) - { - if( otherDevices[ i ] != device ) - { - cl_device_type newType; - error = clGetDeviceInfo( otherDevices[ i ], CL_DEVICE_TYPE, sizeof( newType ), &newType, NULL ); - if( error != CL_SUCCESS ) - { - print_error( error, "Unable to get device type for other device" ); - free(otherDevices); - return NULL; - } - cl_device_id result = otherDevices[ i ]; - free(otherDevices); - return result; - } - } - - // Should never get here - free(otherDevices); - return NULL; -} - -Version get_device_cl_version(cl_device_id device) -{ - size_t str_size; - cl_int err = clGetDeviceInfo(device, CL_DEVICE_VERSION, 0, NULL, &str_size); - ASSERT_SUCCESS(err, "clGetDeviceInfo"); - - std::vector str(str_size); - err = clGetDeviceInfo(device, CL_DEVICE_VERSION, str_size, str.data(), NULL); - ASSERT_SUCCESS(err, "clGetDeviceInfo"); - - if (strstr(str.data(), "OpenCL 1.0") != NULL) - return Version(1, 0); - else if (strstr(str.data(), "OpenCL 1.1") != NULL) - return Version(1, 1); - else if (strstr(str.data(), "OpenCL 1.2") != NULL) - return Version(1, 2); - else if (strstr(str.data(), "OpenCL 2.0") != NULL) - return Version(2, 0); - else if (strstr(str.data(), "OpenCL 2.1") != NULL) - return Version(2, 1); - else if (strstr(str.data(), "OpenCL 2.2") != NULL) - return Version(2, 2); - - throw std::runtime_error(std::string("Unknown OpenCL version: ") + str.data()); -} diff --git a/test_conformance/compatibility/test_common/harness/testHarness.h b/test_conformance/compatibility/test_common/harness/testHarness.h deleted file mode 100644 index 9c2402a9..00000000 --- a/test_conformance/compatibility/test_common/harness/testHarness.h +++ /dev/null @@ -1,147 +0,0 @@ -// -// Copyright (c) 2017-2019 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#ifndef _testHarness_h -#define _testHarness_h - -#include "threadTesting.h" -#include "clImageHelper.h" -#include -#include - -#ifdef __cplusplus -extern "C" { -#endif - -#define ADD_TEST(fn) {test_##fn, #fn, Version(1, 0)} -#define ADD_TEST_VERSION(fn, ver) {test_##fn, #fn, ver} -#define NOT_IMPLEMENTED_TEST(fn) {NULL, #fn, Version(0, 0)} - -#define ARRAY_SIZE(arr) (sizeof(arr) / sizeof((arr)[0])) - -class Version -{ -public: - Version() : m_major(0), m_minor(0) {} - Version(int major, int minor) : m_major(major), m_minor(minor) {} - bool operator>(const Version& rhs) const { return to_int() > rhs.to_int(); } - int to_int() const { return m_major * 10 + m_minor; } - std::string to_string() const - { - std::stringstream ss; - ss << m_major << "." << m_minor; - return ss.str(); - } - -private: - int m_major; - int m_minor; -}; - -typedef struct test_definition -{ - basefn func; - const char* name; - Version min_version; -} test_definition; - -typedef enum test_status -{ - TEST_PASS = 0, - TEST_FAIL = 1, - TEST_SKIP = 2, -} test_status; - -extern int gFailCount; -extern int gTestCount; -extern cl_uint gReSeed; -extern cl_uint gRandomSeed; - -// Supply a list of functions to test here. This will allocate a CL device, create a context, all that -// setup work, and then call each function in turn as dictatated by the passed arguments. -// Returns EXIT_SUCCESS iff all tests succeeded or the tests were listed, -// otherwise return EXIT_FAILURE. -extern int runTestHarness(int argc, const char *argv[], int testNum, test_definition testList[], - int imageSupportRequired, int forceNoContextCreation, cl_command_queue_properties queueProps ); - -// Device checking function. See runTestHarnessWithCheck. If this function returns anything other than TEST_PASS, the harness exits. -typedef test_status (*DeviceCheckFn)( cl_device_id device ); - -// Same as runTestHarness, but also supplies a function that checks the created device for required functionality. -// Returns EXIT_SUCCESS iff all tests succeeded or the tests were listed, -// otherwise return EXIT_FAILURE. -extern int runTestHarnessWithCheck( int argc, const char *argv[], int testNum, test_definition testList[], - int imageSupportRequired, int forceNoContextCreation, cl_command_queue_properties queueProps, - DeviceCheckFn deviceCheckFn ); - -// The command line parser used by runTestHarness to break up parameters into calls to callTestFunctions -extern int parseAndCallCommandLineTests( int argc, const char *argv[], cl_device_id device, int testNum, - test_definition testList[], int forceNoContextCreation, - cl_command_queue_properties queueProps, int num_elements ); - -// Call this function if you need to do all the setup work yourself, and just need the function list called/ -// managed. -// testList is the data structure that contains test functions and its names -// selectedTestList is an array of integers (treated as bools) which tell which function is to be called, -// each element at index i, corresponds to the element in testList at index i -// resultTestList is an array of statuses which contain the result of each selected test -// testNum is the number of tests in testList, selectedTestList and resultTestList -// contextProps are used to create a testing context for each test -// deviceToUse and numElementsToUse are all just passed to each test function -extern void callTestFunctions( test_definition testList[], unsigned char selectedTestList[], test_status resultTestList[], - int testNum, cl_device_id deviceToUse, int forceNoContextCreation, int numElementsToUse, - cl_command_queue_properties queueProps ); - -// This function is called by callTestFunctions, once per function, to do setup, call, logging and cleanup -extern test_status callSingleTestFunction( test_definition test, cl_device_id deviceToUse, int forceNoContextCreation, - int numElementsToUse, cl_command_queue_properties queueProps ); - -///// Miscellaneous steps - -// Given a pre-existing device type choice, check the environment for an override, then print what -// choice was made and how (and return the overridden choice, if there is one) -extern void checkDeviceTypeOverride( cl_device_type *inOutType ); - -// standard callback function for context pfn_notify -extern void CL_CALLBACK notify_callback(const char *errinfo, const void *private_info, size_t cb, void *user_data); - -extern cl_device_type GetDeviceType( cl_device_id ); - -// Given a device (most likely passed in by the harness, but not required), will attempt to find -// a DIFFERENT device and return it. Useful for finding another device to run multi-device tests against. -// Note that returning NULL means an error was hit, but if no error was hit and the device passed in -// is the only device available, the SAME device is returned, so check! -extern cl_device_id GetOpposingDevice( cl_device_id device ); - -Version get_device_cl_version(cl_device_id device); - - -extern int gFlushDenormsToZero; // This is set to 1 if the device does not support denorms (CL_FP_DENORM) -extern int gInfNanSupport; // This is set to 1 if the device supports infinities and NaNs -extern int gIsEmbedded; // This is set to 1 if the device is an embedded device -extern int gHasLong; // This is set to 1 if the device suppots long and ulong types in OpenCL C. -extern int gIsOpenCL_C_1_0_Device; // This is set to 1 if the device supports only OpenCL C 1.0. - -#if ! defined( __APPLE__ ) - void memset_pattern4(void *, const void *, size_t); -#endif - -#ifdef __cplusplus -} -#endif - -#endif // _testHarness_h - - diff --git a/test_conformance/compatibility/test_common/harness/typeWrappers.h b/test_conformance/compatibility/test_common/harness/typeWrappers.h deleted file mode 100644 index 32f8966a..00000000 --- a/test_conformance/compatibility/test_common/harness/typeWrappers.h +++ /dev/null @@ -1,333 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#ifndef _typeWrappers_h -#define _typeWrappers_h - -#include -#include - -#if !defined(_WIN32) -#include -#endif - -#include "compat.h" -#include -#include "mt19937.h" -#include "errorHelpers.h" -#include "kernelHelpers.h" - -extern "C" cl_uint gReSeed; -extern "C" cl_uint gRandomSeed; - -/* cl_context wrapper */ - -class clContextWrapper -{ - public: - clContextWrapper() { mContext = NULL; } - clContextWrapper( cl_context program ) { mContext = program; } - ~clContextWrapper() { if( mContext != NULL ) clReleaseContext( mContext ); } - - clContextWrapper & operator=( const cl_context &rhs ) { mContext = rhs; return *this; } - operator cl_context() { return mContext; } - - cl_context * operator&() { return &mContext; } - - bool operator==( const cl_context &rhs ) { return mContext == rhs; } - - protected: - - cl_context mContext; -}; - -/* cl_program wrapper */ - -class clProgramWrapper -{ - public: - clProgramWrapper() { mProgram = NULL; } - clProgramWrapper( cl_program program ) { mProgram = program; } - ~clProgramWrapper() { if( mProgram != NULL ) clReleaseProgram( mProgram ); } - - clProgramWrapper & operator=( const cl_program &rhs ) { mProgram = rhs; return *this; } - operator cl_program() { return mProgram; } - - cl_program * operator&() { return &mProgram; } - - bool operator==( const cl_program &rhs ) { return mProgram == rhs; } - - protected: - - cl_program mProgram; -}; - -/* cl_kernel wrapper */ - -class clKernelWrapper -{ - public: - clKernelWrapper() { mKernel = NULL; } - clKernelWrapper( cl_kernel kernel ) { mKernel = kernel; } - ~clKernelWrapper() { if( mKernel != NULL ) clReleaseKernel( mKernel ); } - - clKernelWrapper & operator=( const cl_kernel &rhs ) { mKernel = rhs; return *this; } - operator cl_kernel() { return mKernel; } - - cl_kernel * operator&() { return &mKernel; } - - bool operator==( const cl_kernel &rhs ) { return mKernel == rhs; } - - protected: - - cl_kernel mKernel; -}; - -/* cl_mem (stream) wrapper */ - -class clMemWrapper -{ - public: - clMemWrapper() { mMem = NULL; } - clMemWrapper( cl_mem mem ) { mMem = mem; } - ~clMemWrapper() { if( mMem != NULL ) clReleaseMemObject( mMem ); } - - clMemWrapper & operator=( const cl_mem &rhs ) { mMem = rhs; return *this; } - operator cl_mem() { return mMem; } - - cl_mem * operator&() { return &mMem; } - - bool operator==( const cl_mem &rhs ) { return mMem == rhs; } - - protected: - - cl_mem mMem; -}; - -class clProtectedImage -{ - public: - clProtectedImage() { image = NULL; backingStore = NULL; } - clProtectedImage( cl_context context, cl_mem_flags flags, const cl_image_format *fmt, size_t width, cl_int *errcode_ret ); - clProtectedImage( cl_context context, cl_mem_flags flags, const cl_image_format *fmt, size_t width, size_t height, cl_int *errcode_ret ); - clProtectedImage( cl_context context, cl_mem_flags flags, const cl_image_format *fmt, size_t width, size_t height, size_t depth, cl_int *errcode_ret ); - clProtectedImage( cl_context context, cl_mem_object_type imageType, cl_mem_flags flags, const cl_image_format *fmt, size_t width, size_t height, size_t depth, size_t arraySize, cl_int *errcode_ret ); - ~clProtectedImage() - { - if( image != NULL ) - clReleaseMemObject( image ); - -#if defined( __APPLE__ ) - if(backingStore) - munmap(backingStore, backingStoreSize); -#endif - } - - cl_int Create( cl_context context, cl_mem_flags flags, const cl_image_format *fmt, size_t width ); - cl_int Create( cl_context context, cl_mem_flags flags, const cl_image_format *fmt, size_t width, size_t height ); - cl_int Create( cl_context context, cl_mem_flags flags, const cl_image_format *fmt, size_t width, size_t height, size_t depth ); - cl_int Create( cl_context context, cl_mem_object_type imageType, cl_mem_flags flags, const cl_image_format *fmt, size_t width, size_t height, size_t depth, size_t arraySize ); - - clProtectedImage & operator=( const cl_mem &rhs ) { image = rhs; backingStore = NULL; return *this; } - operator cl_mem() { return image; } - - cl_mem * operator&() { return ℑ } - - bool operator==( const cl_mem &rhs ) { return image == rhs; } - - protected: - void *backingStore; - size_t backingStoreSize; - cl_mem image; -}; - -/* cl_command_queue wrapper */ - -class clCommandQueueWrapper -{ - public: - clCommandQueueWrapper() { mMem = NULL; } - clCommandQueueWrapper( cl_command_queue mem ) { mMem = mem; } - ~clCommandQueueWrapper() { if( mMem != NULL ) {int error = clFinish(mMem); if (error) print_error(error, "clFinish failed"); clReleaseCommandQueue( mMem );} } - - clCommandQueueWrapper & operator=( const cl_command_queue &rhs ) { mMem = rhs; return *this; } - operator cl_command_queue() { return mMem; } - - cl_command_queue * operator&() { return &mMem; } - - bool operator==( const cl_command_queue &rhs ) { return mMem == rhs; } - - protected: - - cl_command_queue mMem; -}; - -/* cl_sampler wrapper */ -class clSamplerWrapper -{ - public: - clSamplerWrapper() { mMem = NULL; } - clSamplerWrapper( cl_sampler mem ) { mMem = mem; } - ~clSamplerWrapper() { if( mMem != NULL ) clReleaseSampler( mMem ); } - - clSamplerWrapper & operator=( const cl_sampler &rhs ) { mMem = rhs; return *this; } - operator cl_sampler() { return mMem; } - - cl_sampler * operator&() { return &mMem; } - - bool operator==( const cl_sampler &rhs ) { return mMem == rhs; } - - protected: - - cl_sampler mMem; -}; - -/* cl_event wrapper */ -class clEventWrapper -{ - public: - clEventWrapper() { mMem = NULL; } - clEventWrapper( cl_event mem ) { mMem = mem; } - ~clEventWrapper() { if( mMem != NULL ) clReleaseEvent( mMem ); } - - clEventWrapper & operator=( const cl_event &rhs ) { mMem = rhs; return *this; } - operator cl_event() { return mMem; } - - cl_event * operator&() { return &mMem; } - - bool operator==( const cl_event &rhs ) { return mMem == rhs; } - - protected: - - cl_event mMem; -}; - -/* Generic protected memory buffer, for verifying access within bounds */ -class clProtectedArray -{ - public: - clProtectedArray(); - clProtectedArray( size_t sizeInBytes ); - virtual ~clProtectedArray(); - - void Allocate( size_t sizeInBytes ); - - operator void *() { return (void *)mValidBuffer; } - operator const void *() const { return (const void *)mValidBuffer; } - - protected: - - char * mBuffer; - char * mValidBuffer; - size_t mRealSize, mRoundedSize; -}; - -class RandomSeed -{ - public: - RandomSeed( cl_uint seed ){ if(seed) log_info( "(seed = %10.10u) ", seed ); mtData = init_genrand(seed); } - ~RandomSeed() - { - if( gReSeed ) - gRandomSeed = genrand_int32( mtData ); - free_mtdata(mtData); - } - - operator MTdata () {return mtData;} - - protected: - MTdata mtData; -}; - -template class BufferOwningPtr -{ - BufferOwningPtr(BufferOwningPtr const &); // do not implement - void operator=(BufferOwningPtr const &); // do not implement - - void *ptr; - void *map; - size_t mapsize; // Bytes allocated total, pointed to by map. - size_t allocsize; // Bytes allocated in unprotected pages, pointed to by ptr. - bool aligned; - public: - explicit BufferOwningPtr(void *p = 0) : ptr(p), map(0), mapsize(0), allocsize(0), aligned(false) {} - explicit BufferOwningPtr(void *p, void *m, size_t s) - : ptr(p), map(m), mapsize(s), allocsize(0), aligned(false) - { -#if ! defined( __APPLE__ ) - if(m) - { - log_error( "ERROR: unhandled code path. BufferOwningPtr allocated with mapped buffer!" ); - abort(); - } -#endif - } - ~BufferOwningPtr() { - if (map) { -#if defined( __APPLE__ ) - int error = munmap(map, mapsize); - if (error) log_error("WARNING: munmap failed in BufferOwningPtr.\n"); -#endif - } else { - if ( aligned ) - { - align_free(ptr); - } - else - { - free(ptr); - } - } - } - void reset(void *p, void *m = 0, size_t mapsize_ = 0, size_t allocsize_ = 0, bool aligned_ = false) { - if (map){ -#if defined( __APPLE__ ) - int error = munmap(map, mapsize); - if (error) log_error("WARNING: munmap failed in BufferOwningPtr.\n"); -#else - log_error( "ERROR: unhandled code path. BufferOwningPtr reset with mapped buffer!" ); - abort(); -#endif - } else { - if ( aligned ) - { - align_free(ptr); - } - else - { - free(ptr); - } - } - ptr = p; - map = m; - mapsize = mapsize_; - allocsize = allocsize_; - aligned = aligned_; -#if ! defined( __APPLE__ ) - if(m) - { - log_error( "ERROR: unhandled code path. BufferOwningPtr allocated with mapped buffer!" ); - abort(); - } -#endif - } - operator T*() { return (T*)ptr; } - - size_t getSize() const { return allocsize; }; -}; - -#endif // _typeWrappers_h - - diff --git a/test_conformance/compatibility/test_conformance/CMakeLists.txt b/test_conformance/compatibility/test_conformance/CMakeLists.txt deleted file mode 100644 index 68baa0b9..00000000 --- a/test_conformance/compatibility/test_conformance/CMakeLists.txt +++ /dev/null @@ -1,4 +0,0 @@ - -set(HARNESS_LIB harness-compat) - -add_subdirectory(basic) diff --git a/test_conformance/compatibility/test_conformance/basic/CMakeLists.txt b/test_conformance/compatibility/test_conformance/basic/CMakeLists.txt deleted file mode 100644 index 7caa4ea9..00000000 --- a/test_conformance/compatibility/test_conformance/basic/CMakeLists.txt +++ /dev/null @@ -1,11 +0,0 @@ -set(MODULE_NAME COMPATIBILITY_BASIC) - -set(${MODULE_NAME}_SOURCES - main.c - test_readimage.c - test_writeimage.c -) - -set(${MODULE_NAME}_LIBS harness-compat) - -include(../../../CMakeCommon.txt) diff --git a/test_conformance/compatibility/test_conformance/basic/main.c b/test_conformance/compatibility/test_conformance/basic/main.c deleted file mode 100644 index 6a123dca..00000000 --- a/test_conformance/compatibility/test_conformance/basic/main.c +++ /dev/null @@ -1,44 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "harness/compat.h" - -#if !defined(_WIN32) -#include -#endif - -#include -#include -#include "harness/testHarness.h" -#include "procs.h" - -// FIXME: To use certain functions in harness/imageHelpers.h -// (for example, generate_random_image_data()), the tests are required to declare -// the following variables: -cl_device_type gDeviceType = CL_DEVICE_TYPE_DEFAULT; -bool gTestRounding = false; - -test_definition test_list[] = { - ADD_TEST( readimage ), - ADD_TEST( writeimage ), -}; - -const int test_num = ARRAY_SIZE( test_list ); - -int main(int argc, const char *argv[]) -{ - return runTestHarness( argc, argv, test_num, test_list, false, false, 0 ); -} - diff --git a/test_conformance/compatibility/test_conformance/basic/procs.h b/test_conformance/compatibility/test_conformance/basic/procs.h deleted file mode 100644 index 087e95e7..00000000 --- a/test_conformance/compatibility/test_conformance/basic/procs.h +++ /dev/null @@ -1,142 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "harness/kernelHelpers.h" -#include "harness/testHarness.h" -#include "harness/errorHelpers.h" -#include "harness/typeWrappers.h" -#include "harness/conversions.h" -#include "harness/rounding_mode.h" - -extern void memset_pattern4(void *dest, const void *src_pattern, size_t bytes ); - -extern int test_hostptr(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_fpmath_float(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_fpmath_float2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_fpmath_float4(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_intmath_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_intmath_int2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_intmath_int4(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_intmath_long(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_intmath_long2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_intmath_long4(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_hiloeo(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_if(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_sizeof(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_loop(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_pointer_cast(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_local_arg_def(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_local_kernel_def(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_local_kernel_scope(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_constant(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_constant_source(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_readimage(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_readimage_int16(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_readimage_fp32(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_writeimage(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_writeimage_int16(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_writeimage_fp32(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_mri_one(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_mri_multiple(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_image_r8(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_simplebarrier(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_barrier(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_int2float(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_float2int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_imagearraycopy(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_imagearraycopy3d(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_imagereadwrite(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_imagereadwrite3d(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_readimage3d(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_readimage3d_int16(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_readimage3d_fp32(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_bufferreadwriterect(cl_device_id device, cl_context context, cl_command_queue queue_, int num_elements); -extern int test_imagecopy(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_imagecopy3d(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_imagerandomcopy(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_arraycopy(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems); -extern int test_arrayimagecopy(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_arrayimagecopy3d(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_imagenpot(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_sampler_float(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_sampler_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_createkernelsinprogram(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_single_large_allocation(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_multiple_max_allocation(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_arrayreadwrite(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_imagedim_pow2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_imagedim_non_pow2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_image_param(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_image_multipass_integer_coord(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_image_multipass_float_coord(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_vload_global(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_vload_local(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_vload_constant(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_vload_private(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_vstore_global(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_vstore_local(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_vstore_private(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_explicit_s2v_bool(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_char(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_uchar(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_short(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_ushort(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_int(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_uint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_long(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_ulong(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_float(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_explicit_s2v_double(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_enqueue_map_buffer(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_enqueue_map_image(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_work_item_functions(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_astype(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_native_kernel(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_async_copy_global_to_local(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_async_copy_local_to_global(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_async_strided_copy_global_to_local(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_async_strided_copy_local_to_global(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_prefetch(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_host_numeric_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_kernel_numeric_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_kernel_limit_constants(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_kernel_call_kernel_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - -extern int test_parameter_types(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements); -extern int test_vector_creation(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); -extern int test_vec_type_hint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements); - - -extern int test_kernel_memory_alignment_local(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems ); -extern int test_kernel_memory_alignment_global(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems ); -extern int test_kernel_memory_alignment_constant(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems ); -extern int test_kernel_memory_alignment_private(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems ); - -extern int test_global_work_offsets(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems ); -extern int test_get_global_offset(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems ); - - - diff --git a/test_conformance/compatibility/test_conformance/basic/test_readimage.c b/test_conformance/compatibility/test_conformance/basic/test_readimage.c deleted file mode 100644 index 0888d0e1..00000000 --- a/test_conformance/compatibility/test_conformance/basic/test_readimage.c +++ /dev/null @@ -1,244 +0,0 @@ -// -// Copyright (c) 2017 The Khronos Group Inc. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. -// -#include "harness/compat.h" - -#include -#include -#include -#include - - -#include "procs.h" - -static const char *bgra8888_kernel_code = -"\n" -"__kernel void test_bgra8888(read_only image2d_t srcimg, __global uchar4 *dst, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" int indx = tid_y * get_image_width(srcimg) + tid_x;\n" -" float4 color;\n" -"\n" -" color = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y)) * 255.0f;\n" -" dst[indx] = convert_uchar4_rte(color.zyxw);\n" -"\n" -"}\n"; - - -static const char *rgba8888_kernel_code = -"\n" -"__kernel void test_rgba8888(read_only image2d_t srcimg, __global uchar4 *dst, sampler_t sampler)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" int indx = tid_y * get_image_width(srcimg) + tid_x;\n" -" float4 color;\n" -"\n" -" color = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y)) * 255.0f;\n" -" dst[indx] = convert_uchar4_rte(color);\n" -"\n" -"}\n"; - - -static unsigned char * -generate_8888_image(int w, int h, MTdata d) -{ - unsigned char *ptr = (unsigned char*)malloc(w * h * 4); - int i; - - for (i=0; i -#include -#include -#include - - -#include "procs.h" - -static const char *bgra8888_write_kernel_code = -"\n" -"__kernel void test_bgra8888_write(__global unsigned char *src, write_only image2d_t dstimg)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" int indx = tid_y * get_image_width(dstimg) + tid_x;\n" -" float4 color;\n" -"\n" -" indx *= 4;\n" -" color = (float4)((float)src[indx+2], (float)src[indx+1], (float)src[indx+0], (float)src[indx+3]);\n" -" color /= (float4)(255.0f, 255.0f, 255.0f, 255.0f);\n" -" write_imagef(dstimg, (int2)(tid_x, tid_y), color);\n" -"\n" -"}\n"; - - -static const char *rgba8888_write_kernel_code = -"\n" -"__kernel void test_rgba8888_write(__global unsigned char *src, write_only image2d_t dstimg)\n" -"{\n" -" int tid_x = get_global_id(0);\n" -" int tid_y = get_global_id(1);\n" -" int indx = tid_y * get_image_width(dstimg) + tid_x;\n" -" float4 color;\n" -"\n" -" indx *= 4;\n" -" color = (float4)((float)src[indx+0], (float)src[indx+1], (float)src[indx+2], (float)src[indx+3]);\n" -" color /= (float4)(255.0f, 255.0f, 255.0f, 255.0f);\n" -" write_imagef(dstimg, (int2)(tid_x, tid_y), color);\n" -"\n" -"}\n"; - - -static unsigned char * -generate_8888_image(int w, int h, MTdata d) -{ - cl_uchar *ptr = (cl_uchar *)malloc(w * h * 4); - int i; - - for (i=0; i -#include -#include -#include -#include - -#include "harness/conversions.h" -#include "harness/typeWrappers.h" -#include "harness/kernelHelpers.h" -#include "harness/imageHelpers.h" -#include "harness/errorHelpers.h" -#include "harness/mt19937.h" -#include "harness/rounding_mode.h" -#include "harness/clImageHelper.h" - -extern int gTestCount; -extern int gTestFailure; -extern cl_device_type gDeviceType; - -// Number of iterations per image format to test if not testing max images, rounding, or small images -#define NUM_IMAGE_ITERATIONS 3 - -// Definition for our own sampler type, to mirror the cl_sampler internals -typedef struct { - cl_addressing_mode addressing_mode; - cl_filter_mode filter_mode; - bool normalized_coords; -} image_sampler_data; - -extern void print_read_header( cl_image_format *format, image_sampler_data *sampler, bool err = false, int t = 0 ); -extern void print_write_header( cl_image_format *format, bool err); -extern void print_header( cl_image_format *format, bool err ); -extern bool find_format( cl_image_format *formatList, unsigned int numFormats, cl_image_format *formatToFind ); -extern bool check_minimum_supported( cl_image_format *formatList, unsigned int numFormats, cl_mem_flags flags ); - -cl_channel_type get_channel_type_from_name( const char *name ); -cl_channel_order get_channel_order_from_name( const char *name ); -int random_in_range( int minV, int maxV, MTdata d ); -int random_log_in_range( int minV, int maxV, MTdata d ); - -typedef struct -{ - size_t width; - size_t height; - size_t depth; - size_t rowPitch; - size_t slicePitch; - size_t arraySize; - cl_image_format *format; - cl_mem buffer; - cl_mem_object_type type; -} image_descriptor; - -typedef struct -{ - float p[4]; -}FloatPixel; - -void get_max_sizes(size_t *numberOfSizes, const int maxNumberOfSizes, - size_t sizes[][3], size_t maxWidth, size_t maxHeight, size_t maxDepth, size_t maxArraySize, - const cl_ulong maxIndividualAllocSize, const cl_ulong maxTotalAllocSize, cl_mem_object_type image_type, cl_image_format *format); -extern size_t get_format_max_int( cl_image_format *format ); - -extern char * generate_random_image_data( image_descriptor *imageInfo, BufferOwningPtr &Owner, MTdata d ); - -extern int debug_find_vector_in_image( void *imagePtr, image_descriptor *imageInfo, - void *vectorToFind, size_t vectorSize, int *outX, int *outY, int *outZ ); - -extern int debug_find_pixel_in_image( void *imagePtr, image_descriptor *imageInfo, - unsigned int *valuesToFind, int *outX, int *outY, int *outZ ); -extern int debug_find_pixel_in_image( void *imagePtr, image_descriptor *imageInfo, - int *valuesToFind, int *outX, int *outY, int *outZ ); -extern int debug_find_pixel_in_image( void *imagePtr, image_descriptor *imageInfo, - float *valuesToFind, int *outX, int *outY, int *outZ ); - -extern void copy_image_data( image_descriptor *srcImageInfo, image_descriptor *dstImageInfo, void *imageValues, void *destImageValues, - const size_t sourcePos[], const size_t destPos[], const size_t regionSize[] ); - -int has_alpha(cl_image_format *format); - -inline float calculate_array_index( float coord, float extent ); - -template void read_image_pixel( void *imageData, image_descriptor *imageInfo, - int x, int y, int z, T *outData ) -{ - float convert_half_to_float( unsigned short halfValue ); - - if ( x < 0 || x >= (int)imageInfo->width - || ( imageInfo->height != 0 && ( y < 0 || y >= (int)imageInfo->height ) ) - || ( imageInfo->depth != 0 && ( z < 0 || z >= (int)imageInfo->depth ) ) - || ( imageInfo->arraySize != 0 && ( z < 0 || z >= (int)imageInfo->arraySize ) ) ) - { - // Border color - outData[ 0 ] = outData[ 1 ] = outData[ 2 ] = outData[ 3 ] = 0; - if (!has_alpha(imageInfo->format)) - outData[3] = 1; - return; - } - - cl_image_format *format = imageInfo->format; - - unsigned int i; - T tempData[ 4 ]; - - // Advance to the right spot - char *ptr = (char *)imageData; - size_t pixelSize = get_pixel_size( format ); - - ptr += z * imageInfo->slicePitch + y * imageInfo->rowPitch + x * pixelSize; - - // OpenCL only supports reading floats from certain formats - switch( format->image_channel_data_type ) - { - case CL_SNORM_INT8: - { - cl_char *dPtr = (cl_char *)ptr; - for( i = 0; i < get_format_channel_count( format ); i++ ) - tempData[ i ] = (T)dPtr[ i ]; - break; - } - - case CL_UNORM_INT8: - { - cl_uchar *dPtr = (cl_uchar *)ptr; - for( i = 0; i < get_format_channel_count( format ); i++ ) - tempData[ i ] = (T)dPtr[ i ]; - break; - } - - case CL_SIGNED_INT8: - { - cl_char *dPtr = (cl_char *)ptr; - for( i = 0; i < get_format_channel_count( format ); i++ ) - tempData[ i ] = (T)dPtr[ i ]; - break; - } - - case CL_UNSIGNED_INT8: - { - cl_uchar *dPtr = (cl_uchar*)ptr; - for( i = 0; i < get_format_channel_count( format ); i++ ) - tempData[ i ] = (T)dPtr[ i ]; - break; - } - - case CL_SNORM_INT16: - { - cl_short *dPtr = (cl_short *)ptr; - for( i = 0; i < get_format_channel_count( format ); i++ ) - tempData[ i ] = (T)dPtr[ i ]; - break; - } - - case CL_UNORM_INT16: - { - cl_ushort *dPtr = (cl_ushort *)ptr; - for( i = 0; i < get_format_channel_count( format ); i++ ) - tempData[ i ] = (T)dPtr[ i ]; - break; - } - - case CL_SIGNED_INT16: - { - cl_short *dPtr = (cl_short *)ptr; - for( i = 0; i < get_format_channel_count( format ); i++ ) - tempData[ i ] = (T)dPtr[ i ]; - break; - } - - case CL_UNSIGNED_INT16: - { - cl_ushort *dPtr = (cl_ushort *)ptr; - for( i = 0; i < get_format_channel_count( format ); i++ ) - tempData[ i ] = (T)dPtr[ i ]; - break; - } - - case CL_HALF_FLOAT: - { - cl_ushort *dPtr = (cl_ushort *)ptr; - for( i = 0; i < get_format_channel_count( format ); i++ ) - tempData[ i ] = (T)convert_half_to_float( dPtr[ i ] ); - break; - } - - case CL_SIGNED_INT32: - { - cl_int *dPtr = (cl_int *)ptr; - for( i = 0; i < get_format_channel_count( format ); i++ ) - tempData[ i ] = (T)dPtr[ i ]; - break; - } - - case CL_UNSIGNED_INT32: - { - cl_uint *dPtr = (cl_uint *)ptr; - for( i = 0; i < get_format_channel_count( format ); i++ ) - tempData[ i ] = (T)dPtr[ i ]; - break; - } - - case CL_UNORM_SHORT_565: - { - cl_ushort *dPtr = (cl_ushort*)ptr; - tempData[ 0 ] = (T)( dPtr[ 0 ] >> 11 ); - tempData[ 1 ] = (T)( ( dPtr[ 0 ] >> 5 ) & 63 ); - tempData[ 2 ] = (T)( dPtr[ 0 ] & 31 ); - break; - } - -#ifdef OBSOLETE_FORMAT - case CL_UNORM_SHORT_565_REV: - { - unsigned short *dPtr = (unsigned short *)ptr; - tempData[ 2 ] = (T)( dPtr[ 0 ] >> 11 ); - tempData[ 1 ] = (T)( ( dPtr[ 0 ] >> 5 ) & 63 ); - tempData[ 0 ] = (T)( dPtr[ 0 ] & 31 ); - break; - } - - case CL_UNORM_SHORT_555_REV: - { - unsigned short *dPtr = (unsigned short *)ptr; - tempData[ 2 ] = (T)( ( dPtr[ 0 ] >> 10 ) & 31 ); - tempData[ 1 ] = (T)( ( dPtr[ 0 ] >> 5 ) & 31 ); - tempData[ 0 ] = (T)( dPtr[ 0 ] & 31 ); - break; - } - - case CL_UNORM_INT_8888: - { - unsigned int *dPtr = (unsigned int *)ptr; - tempData[ 3 ] = (T)( dPtr[ 0 ] >> 24 ); - tempData[ 2 ] = (T)( ( dPtr[ 0 ] >> 16 ) & 0xff ); - tempData[ 1 ] = (T)( ( dPtr[ 0 ] >> 8 ) & 0xff ); - tempData[ 0 ] = (T)( dPtr[ 0 ] & 0xff ); - break; - } - case CL_UNORM_INT_8888_REV: - { - unsigned int *dPtr = (unsigned int *)ptr; - tempData[ 0 ] = (T)( dPtr[ 0 ] >> 24 ); - tempData[ 1 ] = (T)( ( dPtr[ 0 ] >> 16 ) & 0xff ); - tempData[ 2 ] = (T)( ( dPtr[ 0 ] >> 8 ) & 0xff ); - tempData[ 3 ] = (T)( dPtr[ 0 ] & 0xff ); - break; - } - - case CL_UNORM_INT_101010_REV: - { - unsigned int *dPtr = (unsigned int *)ptr; - tempData[ 2 ] = (T)( ( dPtr[ 0 ] >> 20 ) & 0x3ff ); - tempData[ 1 ] = (T)( ( dPtr[ 0 ] >> 10 ) & 0x3ff ); - tempData[ 0 ] = (T)( dPtr[ 0 ] & 0x3ff ); - break; - } -#endif - case CL_UNORM_SHORT_555: - { - cl_ushort *dPtr = (cl_ushort *)ptr; - tempData[ 0 ] = (T)( ( dPtr[ 0 ] >> 10 ) & 31 ); - tempData[ 1 ] = (T)( ( dPtr[ 0 ] >> 5 ) & 31 ); - tempData[ 2 ] = (T)( dPtr[ 0 ] & 31 ); - break; - } - - case CL_UNORM_INT_101010: - { - cl_uint *dPtr = (cl_uint *)ptr; - tempData[ 0 ] = (T)( ( dPtr[ 0 ] >> 20 ) & 0x3ff ); - tempData[ 1 ] = (T)( ( dPtr[ 0 ] >> 10 ) & 0x3ff ); - tempData[ 2 ] = (T)( dPtr[ 0 ] & 0x3ff ); - break; - } - - case CL_FLOAT: - { - cl_float *dPtr = (cl_float *)ptr; - for( i = 0; i < get_format_channel_count( format ); i++ ) - tempData[ i ] = (T)dPtr[ i ]; - break; - } -#ifdef CL_SFIXED14_APPLE - case CL_SFIXED14_APPLE: - { - cl_float *dPtr = (cl_float *)ptr; - for( i = 0; i < get_format_channel_count( format ); i++ ) - tempData[ i ] = (T)dPtr[ i ] + 0x4000; - break; - } -#endif - } - - - outData[ 0 ] = outData[ 1 ] = outData[ 2 ] = 0; - outData[ 3 ] = 1; - - if( format->image_channel_order == CL_A ) - { - outData[ 3 ] = tempData[ 0 ]; - } - else if( format->image_channel_order == CL_R ) - { - outData[ 0 ] = tempData[ 0 ]; - } - else if( format->image_channel_order == CL_Rx ) - { - outData[ 0 ] = tempData[ 0 ]; - } - else if( format->image_channel_order == CL_RA ) - { - outData[ 0 ] = tempData[ 0 ]; - outData[ 3 ] = tempData[ 1 ]; - } - else if( format->image_channel_order == CL_RG ) - { - outData[ 0 ] = tempData[ 0 ]; - outData[ 1 ] = tempData[ 1 ]; - } - else if( format->image_channel_order == CL_RGx ) - { - outData[ 0 ] = tempData[ 0 ]; - outData[ 1 ] = tempData[ 1 ]; - } - else if( format->image_channel_order == CL_RGB ) - { - outData[ 0 ] = tempData[ 0 ]; - outData[ 1 ] = tempData[ 1 ]; - outData[ 2 ] = tempData[ 2 ]; - } - else if( format->image_channel_order == CL_RGBx ) - { - outData[ 0 ] = tempData[ 0 ]; - outData[ 1 ] = tempData[ 1 ]; - outData[ 2 ] = tempData[ 2 ]; - } - else if( format->image_channel_order == CL_RGBA ) - { - outData[ 0 ] = tempData[ 0 ]; - outData[ 1 ] = tempData[ 1 ]; - outData[ 2 ] = tempData[ 2 ]; - outData[ 3 ] = tempData[ 3 ]; - } - else if( format->image_channel_order == CL_ARGB ) - { - outData[ 0 ] = tempData[ 1 ]; - outData[ 1 ] = tempData[ 2 ]; - outData[ 2 ] = tempData[ 3 ]; - outData[ 3 ] = tempData[ 0 ]; - } - else if( format->image_channel_order == CL_BGRA ) - { - outData[ 0 ] = tempData[ 2 ]; - outData[ 1 ] = tempData[ 1 ]; - outData[ 2 ] = tempData[ 0 ]; - outData[ 3 ] = tempData[ 3 ]; - } - else if( format->image_channel_order == CL_INTENSITY ) - { - outData[ 1 ] = tempData[ 0 ]; - outData[ 2 ] = tempData[ 0 ]; - outData[ 3 ] = tempData[ 0 ]; - } - else if( format->image_channel_order == CL_LUMINANCE ) - { - outData[ 1 ] = tempData[ 0 ]; - outData[ 2 ] = tempData[ 0 ]; - } -#ifdef CL_1RGB_APPLE - else if( format->image_channel_order == CL_1RGB_APPLE ) - { - outData[ 0 ] = tempData[ 1 ]; - outData[ 1 ] = tempData[ 2 ]; - outData[ 2 ] = tempData[ 3 ]; - outData[ 3 ] = 0xff; - } -#endif -#ifdef CL_BGR1_APPLE - else if( format->image_channel_order == CL_BGR1_APPLE ) - { - outData[ 0 ] = tempData[ 2 ]; - outData[ 1 ] = tempData[ 1 ]; - outData[ 2 ] = tempData[ 0 ]; - outData[ 3 ] = 0xff; - } -#endif - else - { - log_error("Invalid format:"); - print_header(format, true); - } -} - -// Stupid template rules -bool get_integer_coords( float x, float y, float z, - size_t width, size_t height, size_t depth, - image_sampler_data *imageSampler, image_descriptor *imageInfo, - int &outX, int &outY, int &outZ ); -bool get_integer_coords_offset( float x, float y, float z, - float xAddressOffset, float yAddressOffset, float zAddressOffset, - size_t width, size_t height, size_t depth, - image_sampler_data *imageSampler, image_descriptor *imageInfo, - int &outX, int &outY, int &outZ ); - - -template void sample_image_pixel_offset( void *imageData, image_descriptor *imageInfo, - float x, float y, float z, float xAddressOffset, float yAddressOffset, float zAddressOffset, - image_sampler_data *imageSampler, T *outData ) -{ - int iX, iY, iZ; - - float max_w = imageInfo->width; - float max_h; - float max_d; - - switch (imageInfo->type) { - case CL_MEM_OBJECT_IMAGE1D_ARRAY: - max_h = imageInfo->arraySize; - max_d = 0; - break; - case CL_MEM_OBJECT_IMAGE2D_ARRAY: - max_h = imageInfo->height; - max_d = imageInfo->arraySize; - break; - default: - max_h = imageInfo->height; - max_d = imageInfo->depth; - break; - } - - get_integer_coords_offset( x, y, z, xAddressOffset, yAddressOffset, zAddressOffset, max_w, max_h, max_d, imageSampler, imageInfo, iX, iY, iZ ); - - read_image_pixel( imageData, imageInfo, iX, iY, iZ, outData ); -} - - -template void sample_image_pixel( void *imageData, image_descriptor *imageInfo, - float x, float y, float z, image_sampler_data *imageSampler, T *outData ) -{ - return sample_image_pixel_offset(imageData, imageInfo, x, y, z, 0.0f, 0.0f, 0.0f, imageSampler, outData); -} - -FloatPixel sample_image_pixel_float( void *imageData, image_descriptor *imageInfo, - float x, float y, float z, image_sampler_data *imageSampler, float *outData, int verbose, int *containsDenorms ); - -FloatPixel sample_image_pixel_float_offset( void *imageData, image_descriptor *imageInfo, - float x, float y, float z, float xAddressOffset, float yAddressOffset, float zAddressOffset, - image_sampler_data *imageSampler, float *outData, int verbose, int *containsDenorms ); - - -extern void pack_image_pixel( unsigned int *srcVector, const cl_image_format *imageFormat, void *outData ); -extern void pack_image_pixel( int *srcVector, const cl_image_format *imageFormat, void *outData ); -extern void pack_image_pixel( float *srcVector, const cl_image_format *imageFormat, void *outData ); -extern void pack_image_pixel_error( const float *srcVector, const cl_image_format *imageFormat, const void *results, float *errors ); - -extern char *create_random_image_data( ExplicitType dataType, image_descriptor *imageInfo, BufferOwningPtr &P, MTdata d ); - -// deprecated -// extern bool clamp_image_coord( image_sampler_data *imageSampler, float value, size_t max, int &outValue ); - -extern void get_sampler_kernel_code( image_sampler_data *imageSampler, char *outLine ); -extern float get_max_absolute_error( cl_image_format *format, image_sampler_data *sampler); -extern float get_max_relative_error( cl_image_format *format, image_sampler_data *sampler, int is3D, int isLinearFilter ); - - -#define errMax( _x , _y ) ( (_x) != (_x) ? (_x) : (_x) > (_y) ? (_x) : (_y) ) - -static inline cl_uint abs_diff_uint( cl_uint x, cl_uint y ) -{ - return y > x ? y - x : x - y; -} - -static inline cl_uint abs_diff_int( cl_int x, cl_int y ) -{ - return (cl_uint) (y > x ? y - x : x - y); -} - -static inline cl_float relative_error( float test, float expected ) -{ - // 0-0/0 is 0 in this case, not NaN - if( test == 0.0f && expected == 0.0f ) - return 0.0f; - - return (test - expected) / expected; -} - -extern float random_float(float low, float high); - -class CoordWalker -{ -public: - CoordWalker( void * coords, bool useFloats, size_t vecSize ); - ~CoordWalker(); - - cl_float Get( size_t idx, size_t el ); - -protected: - cl_float * mFloatCoords; - cl_int * mIntCoords; - size_t mVecSize; -}; - -extern int DetectFloatToHalfRoundingMode( cl_command_queue ); // Returns CL_SUCCESS on success - -int inline is_half_nan( cl_ushort half ){ return (half & 0x7fff) > 0x7c00; } - -cl_ushort convert_float_to_half( cl_float f ); -cl_float convert_half_to_float( cl_ushort h ); - - -#endif // _image_helpers_h - - diff --git a/test_conformance/opencl_conformance_tests_20_full.csv b/test_conformance/opencl_conformance_tests_20_full.csv index 07ba440b..cfa465bc 100644 --- a/test_conformance/opencl_conformance_tests_20_full.csv +++ b/test_conformance/opencl_conformance_tests_20_full.csv @@ -90,11 +90,6 @@ Math,math_brute_force/bruteforce Integer Ops,integer_ops/test_integer_ops Half Ops,half/Test_half -# ######################################### -# Compatibility with Previous Versions -# ######################################### -Basic 1.2,compatibility/test_conformance/basic/test_basic - ##################################### # OpenCL 2.0 tests ##################################### diff --git a/test_conformance/opencl_conformance_tests_21_full_spirv.csv b/test_conformance/opencl_conformance_tests_21_full_spirv.csv index fdc15334..2cf35c8e 100644 --- a/test_conformance/opencl_conformance_tests_21_full_spirv.csv +++ b/test_conformance/opencl_conformance_tests_21_full_spirv.csv @@ -90,11 +90,6 @@ Math,math_brute_force/bruteforce -offlineCompiler spir_v cache . Integer Ops,integer_ops/test_integer_ops -offlineCompiler spir_v cache . Half Ops,half/Test_half -offlineCompiler spir_v cache . -# ######################################### -# Compatibility with Previous Versions -# ######################################### -Basic 1.2,compatibility/test_conformance/basic/test_basic - ##################################### # OpenCL 2.0 tests ##################################### diff --git a/test_conformance/opencl_conformance_tests_full.csv b/test_conformance/opencl_conformance_tests_full.csv index 387aa141..f155a726 100644 --- a/test_conformance/opencl_conformance_tests_full.csv +++ b/test_conformance/opencl_conformance_tests_full.csv @@ -90,11 +90,6 @@ Math,math_brute_force/bruteforce Integer Ops,integer_ops/test_integer_ops Half Ops,half/test_half -# ######################################### -# Compatibility with Previous Versions -# ######################################### -Basic 1.2,compatibility/test_conformance/basic/test_basic - ##################################### # OpenCL 2.0 tests ##################################### diff --git a/test_conformance/opencl_conformance_tests_full_no_math_or_conversions.csv b/test_conformance/opencl_conformance_tests_full_no_math_or_conversions.csv index ac1427c6..c6dd1d1b 100644 --- a/test_conformance/opencl_conformance_tests_full_no_math_or_conversions.csv +++ b/test_conformance/opencl_conformance_tests_full_no_math_or_conversions.csv @@ -92,11 +92,6 @@ Contractions,contractions/contractions Integer Ops,integer_ops/test_integer_ops Half Ops,half/test_half -# ######################################### -# Compatibility with Previous Versions -# ######################################### -Basic 1.2,compatibility/test_conformance/basic/test_basic - ##################################### # OpenCL 2.0 tests #####################################