ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-19 06:09:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

Remove last compatibility tests (#572)

Browse Source 
They were a duplicate of the non-compatibility version.

The non-compatibility version assumed that BGRA formats aren't
required for implementations that support the embedded profile,
while the compatibility version didn't.

The unified specification doesn't currently document any format
requirements for the embedded profile but the OpenCL 1.2 specification
did and no BGRA format was part of the list. The consensus from the
Working Group is that BGRA formats are not a requirement of the embedded
profile and the specification will be changed to reflect this
(see https://github.com/KhronosGroup/OpenCL-Docs/issues/201).

Closes #494.

Signed-off-by: Kevin Petit <kevin.petit@arm.com>
This commit is contained in:
Kévin Petit
2020-01-31 11:46:36 +00:00
committed by GitHub
parent 695960df2e
commit cac0ab0918
23 changed files with 0 additions and 4658 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
test_conformance/CMakeLists.txt
View File

@@ -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 )

2
test_conformance/compatibility/CMakeLists.txt
View File

@@ -1,2 +0,0 @@
add_subdirectory(test_common)
add_subdirectory(test_conformance)

14
test_conformance/compatibility/test_common/CMakeLists.txt
View File

@@ -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)

253
test_conformance/compatibility/test_common/harness/clImageHelper.h
View File

@@ -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 <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif
#include <stdio.h>
#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

579
test_conformance/compatibility/test_common/harness/errorHelpers.c
View File

@@ -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 <stdio.h>
#include <stdlib.h>
#include <string.h>
#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;
}

249
test_conformance/compatibility/test_common/harness/imageHelpers.cpp
View File

@@ -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;
}

37
test_conformance/compatibility/test_common/harness/imageHelpers.h
View File

@@ -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

684
test_conformance/compatibility/test_common/harness/kernelHelpers.c
View File

@@ -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;
}

129
test_conformance/compatibility/test_common/harness/kernelHelpers.h
View File

@@ -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 <stdio.h>
#include <stdlib.h>
#if defined (__MINGW32__)
#include <malloc.h>
#endif
#include <string.h>
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/opencl.h>
#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

932
test_conformance/compatibility/test_common/harness/testHarness.c
View File

@@ -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 <stdio.h>
#include <stdlib.h>
#if !defined(_WIN32)
#include <stdbool.h>
#endif
#include <string.h>
#include <cassert>
#include <stdexcept>
#include <vector>
#include "threadTesting.h"
#include "errorHelpers.h"
#include "kernelHelpers.h"
#include "fpcontrol.h"
#if !defined(_WIN32)
#include <unistd.h>
#endif
#include <time.h>
#if !defined (__APPLE__)
#include <CL/cl.h>
#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 [<test name>*] [pid<num>] [id<num>] [<device type>]\n", argv[0] );
log_info( "\t<test name>\tOne or more of: (wildcard character '*') (default *)\n");
log_info( "\tpid<num>\tIndicates platform at index <num> should be used (default 0).\n" );
log_info( "\tid<num>\t\tIndicates device at index <num> should be used (default 0).\n" );
log_info( "\t<device_type>\tcpu|gpu|accelerator|<CL_DEVICE_TYPE_*> (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 : "<undefined>" );
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<char> 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());
}

147
test_conformance/compatibility/test_common/harness/testHarness.h
View File

@@ -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 <string>
#include <sstream>
#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

333
test_conformance/compatibility/test_common/harness/typeWrappers.h
View File

@@ -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 <stdio.h>
#include <stdlib.h>
#if !defined(_WIN32)
#include <sys/mman.h>
#endif
#include "compat.h"
#include <stdio.h>
#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 &image; }
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 <typename T> 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

4
test_conformance/compatibility/test_conformance/CMakeLists.txt
View File

@@ -1,4 +0,0 @@
set(HARNESS_LIB harness-compat)
add_subdirectory(basic)

11
test_conformance/compatibility/test_conformance/basic/CMakeLists.txt
View File

@@ -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)

44
test_conformance/compatibility/test_conformance/basic/main.c
View File

@@ -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 <unistd.h>
#endif
#include <stdio.h>
#include <string.h>
#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 );
}

142
test_conformance/compatibility/test_conformance/basic/procs.h
View File

@@ -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 );

244
test_conformance/compatibility/test_conformance/basic/test_readimage.c
View File

@@ -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 <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#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<w*h*4; i++)
ptr[i] = (unsigned char)genrand_int32( d);
return ptr;
}
static int
verify_bgra8888_image(unsigned char *image, unsigned char *outptr, int w, int h)
{
int i;
for (i=0; i<w*h; i++)
{
if (outptr[i] != image[i])
{
log_error("READ_IMAGE_BGRA_UNORM_INT8 test failed\n");
return -1;
}
}
log_info("READ_IMAGE_BGRA_UNORM_INT8 test passed\n");
return 0;
}
static int
verify_rgba8888_image(unsigned char *image, unsigned char *outptr, int w, int h)
{
int i;
for (i=0; i<w*h*4; i++)
{
if (outptr[i] != image[i])
{
log_error("READ_IMAGE_RGBA_UNORM_INT8 test failed\n");
return -1;
}
}
log_info("READ_IMAGE_RGBA_UNORM_INT8 test passed\n");
return 0;
}
int test_readimage(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
cl_mem streams[3];
cl_program program[2];
cl_kernel kernel[2];
cl_image_format img_format;
unsigned char *input_ptr[2], *output_ptr;
size_t threads[2];
int img_width = 512;
int img_height = 512;
int i, err;
size_t origin[3] = {0, 0, 0};
size_t region[3] = {img_width, img_height, 1};
size_t length = img_width * img_height * 4 * sizeof(unsigned char);
MTdata d;
PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
d = init_genrand( gRandomSeed );
input_ptr[0] = generate_8888_image(img_width, img_height, d);
input_ptr[1] = generate_8888_image(img_width, img_height, d);
free_mtdata(d); d = NULL;
output_ptr = (unsigned char*)malloc(length);
img_format.image_channel_order = CL_BGRA;
img_format.image_channel_data_type = CL_UNORM_INT8;
streams[0] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format, img_width, img_height, 0, NULL, NULL);
if (!streams[0])
{
log_error("create_image_2d failed\n");
return -1;
}
img_format.image_channel_order = CL_RGBA;
img_format.image_channel_data_type = CL_UNORM_INT8;
streams[1] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format, img_width, img_height, 0, NULL, NULL);
if (!streams[1])
{
log_error("create_image_2d failed\n");
return -1;
}
streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
if (!streams[2])
{
log_error("clCreateBuffer failed\n");
return -1;
}
err = clEnqueueWriteImage(queue, streams[0], CL_TRUE, origin, region, 0, 0, input_ptr[0], 0, NULL, NULL);
if (err != CL_SUCCESS)
{
log_error("clEnqueueWriteImage failed\n");
return -1;
}
err = clEnqueueWriteImage(queue, streams[1], CL_TRUE, origin, region, 0, 0, input_ptr[1], 0, NULL, NULL);
if (err != CL_SUCCESS)
{
log_error("clEnqueueWriteImage failed\n");
return -1;
}
err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &bgra8888_kernel_code, "test_bgra8888" );
if (err)
return -1;
err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &rgba8888_kernel_code, "test_rgba8888" );
if (err)
return -1;
cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err);
test_error(err, "clCreateSampler failed");
err = clSetKernelArg(kernel[0], 0, sizeof streams[0], &streams[0]);
err |= clSetKernelArg(kernel[0], 1, sizeof streams[2], &streams[2]);
err |= clSetKernelArg(kernel[0], 2, sizeof sampler, &sampler);
if (err != CL_SUCCESS)
{
log_error("clSetKernelArg failed\n");
return -1;
}
err = clSetKernelArg(kernel[1], 0, sizeof streams[1], &streams[1]);
err |= clSetKernelArg(kernel[1], 1, sizeof streams[2], &streams[2]);
err |= clSetKernelArg(kernel[1], 2, sizeof sampler, &sampler);
if (err != CL_SUCCESS)
{
log_error("clSetKernelArg failed\n");
return -1;
}
threads[0] = (unsigned int)img_width;
threads[1] = (unsigned int)img_height;
for (i=0; i<2; i++)
{
err = clEnqueueNDRangeKernel(queue, kernel[i], 2, NULL, threads, NULL, 0, NULL, NULL);
if (err != CL_SUCCESS)
{
log_error("%s clEnqueueNDRangeKernel failed\n", __FUNCTION__);
return -1;
}
err = clEnqueueReadBuffer(queue, streams[2], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
if (err != CL_SUCCESS)
{
log_error("clEnqueueReadBuffer failed\n");
return -1;
}
switch (i)
{
case 0:
err = verify_bgra8888_image(input_ptr[i], output_ptr, img_width, img_height);
break;
case 1:
err = verify_rgba8888_image(input_ptr[i], output_ptr, img_width, img_height);
break;
}
if (err)
break;
}
// cleanup
clReleaseSampler(sampler);
clReleaseMemObject(streams[0]);
clReleaseMemObject(streams[1]);
clReleaseMemObject(streams[2]);
for (i=0; i<2; i++)
{
clReleaseKernel(kernel[i]);
clReleaseProgram(program[i]);
}
free(input_ptr[0]);
free(input_ptr[1]);
free(output_ptr);
return err;
}

300
test_conformance/compatibility/test_conformance/basic/test_writeimage.c
View File

@@ -1,300 +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 <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#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<w*h*4; i++)
ptr[i] = (cl_uchar)genrand_int32(d);
return ptr;
}
static int
verify_bgra8888_image(unsigned char *image, unsigned char *outptr, int w, int h)
{
int i;
for (i=0; i<w*h*4; i++)
{
if (outptr[i] != image[i])
{
log_error("WRITE_IMAGE_BGRA_UNORM_INT8 test failed\n");
return -1;
}
}
log_info("WRITE_IMAGE_BGRA_UNORM_INT8 test passed\n");
return 0;
}
static int
verify_rgba8888_image(unsigned char *image, unsigned char *outptr, int w, int h)
{
int i;
for (i=0; i<w*h*4; i++)
{
if (outptr[i] != image[i])
{
log_error("WRITE_IMAGE_RGBA_UNORM_INT8 test failed\n");
return -1;
}
}
log_info("WRITE_IMAGE_RGBA_UNORM_INT8 test passed\n");
return 0;
}
int test_writeimage(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
cl_mem streams[6];
cl_program program[2];
cl_kernel kernel[4];
unsigned char *input_ptr[2], *output_ptr;
cl_image_format img_format;
size_t threads[2];
int img_width = 512;
int img_height = 512;
int i, err, any_err = 0;
size_t origin[3] = {0, 0, 0};
size_t region[3] = {img_width, img_height, 1};
size_t length = img_width * img_height * 4 * sizeof(unsigned char);
PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
MTdata d = init_genrand( gRandomSeed );
input_ptr[0] = generate_8888_image(img_width, img_height, d);
input_ptr[1] = generate_8888_image(img_width, img_height, d);
free_mtdata(d); d = NULL;
output_ptr = (unsigned char*)malloc(length);
img_format.image_channel_order = CL_BGRA;
img_format.image_channel_data_type = CL_UNORM_INT8;
streams[0] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format, img_width, img_height, 0, NULL, NULL);
if (!streams[0])
{
log_error("create_image_2d failed\n");
return -1;
}
img_format.image_channel_order = CL_RGBA;
img_format.image_channel_data_type = CL_UNORM_INT8;
streams[1] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format, img_width, img_height, 0, NULL, NULL);
if (!streams[1])
{
log_error("create_image_2d failed\n");
return -1;
}
img_format.image_channel_order = CL_BGRA;
img_format.image_channel_data_type = CL_UNORM_INT8;
streams[2] = create_image_2d(context, CL_MEM_WRITE_ONLY, &img_format, img_width, img_height, 0, NULL, NULL);
if (!streams[2])
{
log_error("create_image_2d failed\n");
return -1;
}
img_format.image_channel_order = CL_RGBA;
img_format.image_channel_data_type = CL_UNORM_INT8;
streams[3] = create_image_2d(context, CL_MEM_WRITE_ONLY, &img_format, img_width, img_height, 0, NULL, NULL);
if (!streams[3])
{
log_error("create_image_2d failed\n");
return -1;
}
streams[4] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
if (!streams[4])
{
log_error("clCreateBuffer failed\n");
return -1;
}
streams[5] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
if (!streams[5])
{
log_error("clCreateBuffer failed\n");
return -1;
}
err = clEnqueueWriteBuffer(queue, streams[4], CL_TRUE, 0, length, input_ptr[0], 0, NULL, NULL);
if (err != CL_SUCCESS)
{
log_error("clEnqueueWriteBuffer failed\n");
return -1;
}
err = clEnqueueWriteBuffer(queue, streams[5], CL_TRUE, 0, length, input_ptr[1], 0, NULL, NULL);
if (err != CL_SUCCESS)
{
log_error("clEnqueueWriteBuffer failed\n");
return -1;
}
err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &bgra8888_write_kernel_code, "test_bgra8888_write" );
if (err)
return -1;
kernel[2] = clCreateKernel(program[0], "test_bgra8888_write", NULL);
if (!kernel[2])
{
log_error("clCreateKernel failed\n");
return -1;
}
err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &rgba8888_write_kernel_code, "test_rgba8888_write" );
if (err)
return -1;
kernel[3] = clCreateKernel(program[1], "test_rgba8888_write", NULL);
if (!kernel[3])
{
log_error("clCreateKernel failed\n");
return -1;
}
err = clSetKernelArg(kernel[0], 0, sizeof streams[4], &streams[4]);
err |= clSetKernelArg(kernel[0], 1, sizeof streams[0], &streams[0]);
if (err != CL_SUCCESS)
{
log_error("clSetKernelArgs failed\n");
return -1;
}
err = clSetKernelArg(kernel[1], 0, sizeof streams[5], &streams[5]);
err |= clSetKernelArg(kernel[1], 1, sizeof streams[1], &streams[1]);
if (err != CL_SUCCESS)
{
log_error("clSetKernelArgs failed\n");
return -1;
}
err = clSetKernelArg(kernel[2], 0, sizeof streams[4], &streams[4]);
err |= clSetKernelArg(kernel[2], 1, sizeof streams[2], &streams[2]);
if (err != CL_SUCCESS)
{
log_error("clSetKernelArgs failed\n");
return -1;
}
err = clSetKernelArg(kernel[3], 0, sizeof streams[5], &streams[5]);
err |= clSetKernelArg(kernel[3], 1, sizeof streams[3], &streams[3]);
if (err != CL_SUCCESS)
{
log_error("clSetKernelArgs failed\n");
return -1;
}
threads[0] = (unsigned int)img_width;
threads[1] = (unsigned int)img_height;
for (i=0; i<4; i++)
{
err = clEnqueueNDRangeKernel(queue, kernel[i], 2, NULL, threads, NULL, 0, NULL, NULL);
if (err != CL_SUCCESS)
{
log_error("clEnqueueNDRangeKernel failed\n");
return -1;
}
err = clEnqueueReadImage(queue, streams[i], CL_TRUE, origin, region, 0, 0, output_ptr, 0, NULL, NULL);
if (err != CL_SUCCESS)
{
log_error("clReadImage failed\n");
return -1;
}
switch (i)
{
case 0:
case 2:
err = verify_bgra8888_image(input_ptr[i&0x01], output_ptr, img_width, img_height);
break;
case 1:
case 3:
err = verify_rgba8888_image(input_ptr[i&0x01], output_ptr, img_width, img_height);
break;
}
//if (err)
//break;
any_err |= err;
}
// cleanup
clReleaseMemObject(streams[0]);
clReleaseMemObject(streams[1]);
clReleaseMemObject(streams[2]);
clReleaseMemObject(streams[3]);
clReleaseMemObject(streams[4]);
clReleaseMemObject(streams[5]);
for (i=0; i<2; i++)
{
clReleaseKernel(kernel[i]);
clReleaseKernel(kernel[i+2]);
clReleaseProgram(program[i]);
}
free(input_ptr[0]);
free(input_ptr[1]);
free(output_ptr);
return any_err;
}

533
test_conformance/compatibility/test_conformance/images/image_helpers.h
View File

@@ -1,533 +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 _image_helpers_h
#define _image_helpers_h
#include "harness/compat.h"
#include <stdio.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#include <time.h>
#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<char> &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 <class T> 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 <class T> 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<T>( imageData, imageInfo, iX, iY, iZ, outData );
}
template <class T> 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<T>(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<char> &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

5
test_conformance/opencl_conformance_tests_20_full.csv
View File

@@ -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
#####################################
1 #
90 C11 Atomics,c11_atomics/test_c11_atomics Pipes,pipes/test_pipes
91 Execution Model,device_execution/test_device_execution SVM,SVM/test_SVM
92 Generic Address Space,generic_address_space/test_generic_address_space Workgroups,workgroups/test_workgroups
Non Uniform Work Groups,non_uniform_work_group/test_non_uniform_work_group
Pipes,pipes/test_pipes
SVM,SVM/test_SVM
Workgroups,workgroups/test_workgroups
#########################################
93 # Extensions #########################################
94 ######################################### # Extensions
95 SPIR,spir/test_spir #########################################

5
test_conformance/opencl_conformance_tests_21_full_spirv.csv
View File

@@ -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
#####################################
1 #
90 C11 Atomics,c11_atomics/test_c11_atomics -offlineCompiler spir_v cache . Pipes,pipes/test_pipes -offlineCompiler spir_v cache .
91 Execution Model,device_execution/test_device_execution -offlineCompiler spir_v cache . SVM,SVM/test_SVM -offlineCompiler spir_v cache .
92 Generic Address Space,generic_address_space/test_generic_address_space -offlineCompiler spir_v cache . Workgroups,workgroups/test_workgroups -offlineCompiler spir_v cache .
Non Uniform Work Groups,non_uniform_work_group/test_non_uniform_work_group -offlineCompiler spir_v cache .
Pipes,pipes/test_pipes -offlineCompiler spir_v cache .
SVM,SVM/test_SVM -offlineCompiler spir_v cache .
Workgroups,workgroups/test_workgroups -offlineCompiler spir_v cache .
#####################################
93 # OpenCL 2.1 tests #####################################
94 ##################################### # OpenCL 2.1 tests
95 Device timer,device_timer/test_device_timer #####################################

5
test_conformance/opencl_conformance_tests_full.csv
View File

@@ -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
#####################################
1 #
90 C11 Atomics,c11_atomics/test_c11_atomics Pipes,pipes/test_pipes
91 Execution Model,device_execution/test_device_execution SVM,SVM/test_SVM
92 Generic Address Space,generic_address_space/test_generic_address_space Workgroups,workgroups/test_workgroups
Non Uniform Work Groups,non_uniform_work_group/test_non_uniform_work_group
Pipes,pipes/test_pipes
SVM,SVM/test_SVM
Workgroups,workgroups/test_workgroups
#####################################
93 # OpenCL 2.1 tests #####################################
94 ##################################### # OpenCL 2.1 tests
95 Device timer,device_timer/test_device_timer #####################################

5
test_conformance/opencl_conformance_tests_full_no_math_or_conversions.csv
View File

@@ -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
#####################################
1 #
92 C11 Atomics,c11_atomics/test_c11_atomics Pipes,pipes/test_pipes
93 Execution Model,device_execution/test_device_execution SVM,SVM/test_SVM
94 Generic Address Space,generic_address_space/test_generic_address_space Workgroups,workgroups/test_workgroups
Non Uniform Work Groups,non_uniform_work_group/test_non_uniform_work_group
Pipes,pipes/test_pipes
SVM,SVM/test_SVM
Workgroups,workgroups/test_workgroups
#####################################
95 # OpenCL 2.1 tests #####################################
96 ##################################### # OpenCL 2.1 tests
97 Device timer,device_timer/test_device_timer #####################################