//
// 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 "harness/rounding_mode.h"
#include "harness/ThreadPool.h"
#include "harness/testHarness.h"
#include "harness/kernelHelpers.h"
#include "harness/parseParameters.h"
#if defined(__APPLE__)
#include <sys/sysctl.h>
#endif
#if defined( __linux__ )
#include <unistd.h>
#include <sys/syscall.h>
#include <linux/sysctl.h>
#endif
#if defined(__linux__)
#include <sys/param.h>
#include <libgen.h>
#endif
#include "mingw_compat.h"
#if defined(__MINGW32__)
#include <sys/param.h>
#endif
#include <stdarg.h>
#include <stdio.h>
#include <string.h>
#if !defined(_WIN32)
#include <libgen.h>
#include <sys/mman.h>
#endif
#include <time.h>
#include "Sleep.h"
#include "basic_test_conversions.h"
#if (defined(_WIN32) && defined (_MSC_VER))
// need for _controlfp_s and rouinding modes in RoundingMode
#include "harness/testHarness.h"
#endif
#pragma mark -
#pragma mark globals
#define BUFFER_SIZE (1024*1024)
#define kPageSize 4096
#define EMBEDDED_REDUCTION_FACTOR 16
#define PERF_LOOP_COUNT 100
#define kCallStyleCount (kVectorSizeCount + 1 /* for implicit scalar */)
#if (defined(__arm__) || defined(__aarch64__)) && defined(__GNUC__)
#include "fplib.h"
extern bool qcom_sat;
extern roundingMode qcom_rm;
#endif
const char ** argList = NULL;
int argCount = 0;
cl_context gContext = NULL;
cl_command_queue gQueue = NULL;
char appName[64] = "ctest";
int gStartTestNumber = -1;
int gEndTestNumber = 0;
#if defined( __APPLE__ )
int gTimeResults = 1;
#else
int gTimeResults = 0;
#endif
int gReportAverageTimes = 0;
void *gIn = NULL;
void *gRef = NULL;
void *gAllowZ = NULL;
void *gOut[ kCallStyleCount ] = { NULL };
cl_mem gInBuffer;
cl_mem gOutBuffers[ kCallStyleCount ];
size_t gComputeDevices = 0;
uint32_t gDeviceFrequency = 0;
int gWimpyMode = 0;
int gWimpyReductionFactor = 128;
int gSkipTesting = 0;
int gForceFTZ = 0;
int gMultithread = 1;
int gIsRTZ = 0;
uint32_t gSimdSize = 1;
int gHasDouble = 0;
int gTestDouble = 1;
const char * sizeNames[] = { "", "", "2", "3", "4", "8", "16" };
const int vectorSizes[] = { 1, 1, 2, 3, 4, 8, 16 };
int gMinVectorSize = 0;
int gMaxVectorSize = sizeof(vectorSizes) / sizeof( vectorSizes[0] );
static MTdata gMTdata;
#pragma mark -
#pragma mark Declarations
static int ParseArgs( int argc, const char **argv );
static void PrintUsage( void );
test_status InitCL( cl_device_id device );
static int GetTestCase( const char *name, Type *outType, Type *inType, SaturationMode *sat, RoundingMode *round );
static int DoTest( cl_device_id device, Type outType, Type inType, SaturationMode sat, RoundingMode round, MTdata d );
static cl_program MakeProgram( Type outType, Type inType, SaturationMode sat, RoundingMode round, int vectorSize, cl_kernel *outKernel );
static int RunKernel( cl_kernel kernel, void *inBuf, void *outBuf, size_t blockCount );
void *FlushToZero( void );
void UnFlushToZero( void *);
static cl_program CreateImplicitConvertProgram( Type outType, Type inType, SaturationMode sat, RoundingMode round, int vectorSize, char testName[256], cl_int *error );
static cl_program CreateStandardProgram( Type outType, Type inType, SaturationMode sat, RoundingMode round, int vectorSize, char testName[256], cl_int *error );
// Windows (since long double got deprecated) sets the x87 to 53-bit precision
// (that's x87 default state). This causes problems with the tests that
// convert long and ulong to float and double or otherwise deal with values
// that need more precision than 53-bit. So, set the x87 to 64-bit precision.
static inline void Force64BitFPUPrecision(void)
{
#if __MINGW32__
// The usual method is to use _controlfp as follows:
// #include <float.h>
// _controlfp(_PC_64, _MCW_PC);
//
// _controlfp is available on MinGW32 but not on MinGW64. Instead of having
// divergent code just use inline assembly which works for both.
unsigned short int orig_cw = 0;
unsigned short int new_cw = 0;
__asm__ __volatile__ ("fstcw %0":"=m" (orig_cw));
new_cw = orig_cw | 0x0300; // set precision to 64-bit
__asm__ __volatile__ ("fldcw %0"::"m" (new_cw));
#else
/* Implement for other platforms if needed */
#endif
}
int test_conversions( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
{
int error, i, testNumber = -1;
int startMinVectorSize = gMinVectorSize;
Type inType, outType;
RoundingMode round;
SaturationMode sat;
if( argCount )
{
for( i = 0; i < argCount; i++ )
{
if( GetTestCase( argList[i], &outType, &inType, &sat, &round ) )
{
vlog_error( "\n\t\t**** ERROR: Unable to parse function name %s. Skipping.... *****\n\n", argList[i] );
continue;
}
// skip double if we don't have it
if( !gTestDouble && (inType == kdouble || outType == kdouble ) )
{
if( gHasDouble )
{
vlog_error( "\t *** convert_%sn%s%s( %sn ) FAILED ** \n", gTypeNames[ outType ], gSaturationNames[ sat ], gRoundingModeNames[round], gTypeNames[inType] );
vlog( "\t\tcl_khr_fp64 enabled, but double testing turned off.\n" );
}
continue;
}
// skip longs on embedded
if( !gHasLong && (inType == klong || outType == klong || inType == kulong || outType == kulong) )
{
continue;
}
// Skip the implicit converts if the rounding mode is not default or test is saturated
if( 0 == startMinVectorSize )
{
if( sat || round != kDefaultRoundingMode )
gMinVectorSize = 1;
else
gMinVectorSize = 0;
}
if( ( error = DoTest( device, outType, inType, sat, round, gMTdata ) ) )
{
vlog_error( "\t *** convert_%sn%s%s( %sn ) FAILED ** \n", gTypeNames[outType], gSaturationNames[sat], gRoundingModeNames[round], gTypeNames[inType] );
}
}
}
else
{
for( outType = (Type)0; outType < kTypeCount; outType = (Type)(outType+1) )
{
for( inType = (Type)0; inType < kTypeCount; inType = (Type)(inType+1) )
{
// skip longs on embedded
if( !gHasLong && (inType == klong || outType == klong || inType == kulong || outType == kulong) )
{
continue;
}
for( sat = (SaturationMode)0; sat < kSaturationModeCount; sat = (SaturationMode)(sat+1) )
{
//skip illegal saturated conversions to float type
if( kSaturated == sat && ( outType == kfloat || outType == kdouble ) )
{
continue;
}
for( round = (RoundingMode)0; round < kRoundingModeCount; round = (RoundingMode)(round+1) )
{
if( ++testNumber < gStartTestNumber )
{
// vlog( "%d) skipping convert_%sn%s%s( %sn )\n", testNumber, gTypeNames[ outType ], gSaturationNames[ sat ], gRoundingModeNames[round], gTypeNames[inType] );
continue;
}
else
{
if( gEndTestNumber > 0 && testNumber >= gEndTestNumber )
{
goto exit;
}
}
vlog( "%d) Testing convert_%sn%s%s( %sn ):\n", testNumber, gTypeNames[ outType ], gSaturationNames[ sat ], gRoundingModeNames[round], gTypeNames[inType] );
// skip double if we don't have it
if( ! gTestDouble && (inType == kdouble || outType == kdouble ) )
{
if( gHasDouble )
{
vlog_error( "\t *** %d) convert_%sn%s%s( %sn ) FAILED ** \n", testNumber, gTypeNames[ outType ], gSaturationNames[ sat ], gRoundingModeNames[round], gTypeNames[inType] );
vlog( "\t\tcl_khr_fp64 enabled, but double testing turned off.\n" );
}
continue;
}
// Skip the implicit converts if the rounding mode is not default or test is saturated
if( 0 == startMinVectorSize )
{
if( sat || round != kDefaultRoundingMode )
gMinVectorSize = 1;
else
gMinVectorSize = 0;
}
if( ( error = DoTest( device, outType, inType, sat, round, gMTdata ) ) )
{
vlog_error( "\t *** %d) convert_%sn%s%s( %sn ) FAILED ** \n", testNumber, gTypeNames[outType], gSaturationNames[sat], gRoundingModeNames[round], gTypeNames[inType] );
}
}
}
}
}
}
exit:
return gFailCount;
}
test_definition test_list[] = {
ADD_TEST( conversions ),
};
const int test_num = ARRAY_SIZE( test_list );
#pragma mark -
int main (int argc, const char **argv )
{
int error;
cl_uint seed = (cl_uint) time( NULL );
argc = parseCustomParam(argc, argv);
if (argc == -1)
{
return 1;
}
if( (error = ParseArgs( argc, argv )) )
return error;
//Turn off sleep so our tests run to completion
PreventSleep();
atexit( ResumeSleep );
if(!gMultithread)
SetThreadCount(1);
#if defined(_MSC_VER) && defined(_M_IX86)
// VS2005 (and probably others, since long double got deprecated) sets
// the x87 to 53-bit precision. This causes problems with the tests
// that convert long and ulong to float and double, since they deal
// with values that need more precision than that. So, set the x87
// to 64-bit precision.
unsigned int ignored;
_controlfp_s(&ignored, _PC_64, _MCW_PC);
#endif
vlog( "===========================================================\n" );
vlog( "Random seed: %u\n", seed );
gMTdata = init_genrand( seed );
const char* arg[] = {argv[0]};
int ret = runTestHarnessWithCheck( 1, arg, test_num, test_list, true, 0, InitCL );
free_mtdata( gMTdata );
if (gQueue)
{
error = clFinish(gQueue);
if (error) vlog_error("clFinish failed: %d\n", error);
}
clReleaseMemObject(gInBuffer);
for( int i = 0; i < kCallStyleCount; i++ ) {
clReleaseMemObject(gOutBuffers[i]);
}
clReleaseCommandQueue(gQueue);
clReleaseContext(gContext);
return ret;
}
#pragma mark -
#pragma mark setup
static int ParseArgs( int argc, const char **argv )
{
int i;
argList = (const char **)calloc( argc - 1, sizeof( char*) );
argCount = 0;
if( NULL == argList && argc > 1 )
return -1;
#if (defined( __APPLE__ ) || defined(__linux__) || defined (__MINGW32__))
{ // Extract the app name
char baseName[ MAXPATHLEN ];
strncpy( baseName, argv[0], MAXPATHLEN );
char *base = basename( baseName );
if( NULL != base )
{
strncpy( appName, base, sizeof( appName ) );
appName[ sizeof( appName ) -1 ] = '\0';
}
}
#elif defined (_WIN32)
{
char fname[_MAX_FNAME + _MAX_EXT + 1];
char ext[_MAX_EXT];
errno_t err = _splitpath_s( argv[0], NULL, 0, NULL, 0,
fname, _MAX_FNAME, ext, _MAX_EXT );
if (err == 0) { // no error
strcat (fname, ext); //just cat them, size of frame can keep both
strncpy (appName, fname, sizeof(appName));
appName[ sizeof( appName ) -1 ] = '\0';
}
}
#endif
vlog( "\n%s", appName );
for( i = 1; i < argc; i++ )
{
const char *arg = argv[i];
if( NULL == arg )
break;
vlog( "\t%s", arg );
if( arg[0] == '-' )
{
arg++;
while( *arg != '\0' )
{
switch( *arg )
{
case 'd':
gTestDouble ^= 1;
break;
case 'l':
gSkipTesting ^= 1;
break;
case 'm':
gMultithread ^= 1;
break;
case 'w':
gWimpyMode ^= 1;
break;
case '[':
parseWimpyReductionFactor(arg, gWimpyReductionFactor);
break;
case 'z':
gForceFTZ ^= 1;
break;
case 't':
gTimeResults ^= 1;
break;
case 'a':
gReportAverageTimes ^= 1;
break;
case '1':
if( arg[1] == '6' )
{
gMinVectorSize = 6;
gMaxVectorSize = 7;
arg++;
}
else
{
gMinVectorSize = 0;
gMaxVectorSize = 2;
}
break;
case '2':
gMinVectorSize = 2;
gMaxVectorSize = 3;
break;
case '3':
gMinVectorSize = 3;
gMaxVectorSize = 4;
break;
case '4':
gMinVectorSize = 4;
gMaxVectorSize = 5;
break;
case '8':
gMinVectorSize = 5;
gMaxVectorSize = 6;
break;
default:
vlog( " <-- unknown flag: %c (0x%2.2x)\n)", *arg, *arg );
PrintUsage();
return -1;
}
arg++;
}
}
else
{
char *t = NULL;
long number = strtol( arg, &t, 0 );
if( t != arg )
{
if( gStartTestNumber != -1 )
gEndTestNumber = gStartTestNumber + (int) number;
else
gStartTestNumber = (int) number;
}
else
{
argList[ argCount ] = arg;
argCount++;
}
}
}
// Check for the wimpy mode environment variable
if (getenv("CL_WIMPY_MODE")) {
vlog( "\n" );
vlog( "*** Detected CL_WIMPY_MODE env ***\n" );
gWimpyMode = 1;
}
vlog( "\n" );
vlog( "Test binary built %s %s\n", __DATE__, __TIME__ );
PrintArch();
if( gWimpyMode )
{
vlog( "\n" );
vlog( "*** WARNING: Testing in Wimpy mode! ***\n" );
vlog( "*** Wimpy mode is not sufficient to verify correctness. ***\n" );
vlog( "*** It gives warm fuzzy feelings and then nevers calls. ***\n\n" );
vlog("*** Wimpy Reduction Factor: %-27u ***\n\n", gWimpyReductionFactor);
}
return 0;
}
static void PrintUsage( void )
{
int i;
vlog( "%s [-wz#]: <optional: test names>\n", appName );
vlog( "\ttest names:\n" );
vlog( "\t\tdestFormat<_sat><_round>_sourceFormat\n" );
vlog( "\t\t\tPossible format types are:\n\t\t\t\t" );
for( i = 0; i < kTypeCount; i++ )
vlog( "%s, ", gTypeNames[i] );
vlog( "\n\n\t\t\tPossible saturation values are: (empty) and _sat\n" );
vlog( "\t\t\tPossible rounding values are:\n\t\t\t\t(empty), " );
for( i = 1; i < kRoundingModeCount; i++ )
vlog( "%s, ", gRoundingModeNames[i] );
vlog( "\n\t\t\tExamples:\n" );
vlog( "\t\t\t\tulong_short converts short to ulong\n" );
vlog( "\t\t\t\tchar_sat_rte_float converts float to char with saturated clipping in round to nearest rounding mode\n\n" );
vlog( "\toptions:\n" );
vlog( "\t\t-d\tToggle testing of double precision. On by default if cl_khr_fp64 is enabled, ignored otherwise.\n" );
vlog( "\t\t-l\tToggle link check mode. When on, testing is skipped, and we just check to see that the kernels build. (Off by default.)\n" );
vlog( "\t\t-m\tToggle Multithreading. (On by default.)\n" );
vlog( "\t\t-w\tToggle wimpy mode. When wimpy mode is on, we run a very small subset of the tests for each fn. NOT A VALID TEST! (Off by default.)\n" );
vlog(" \t\t-[2^n]\tSet wimpy reduction factor, recommended range of n is 1-12, default factor(%u)\n", gWimpyReductionFactor);
vlog( "\t\t-z\tToggle flush to zero mode (Default: per device)\n" );
vlog( "\t\t-#\tTest just vector size given by #, where # is an element of the set {1,2,3,4,8,16}\n" );
vlog( "\n" );
vlog( "You may also pass the number of the test on which to start.\nA second number can be then passed to indicate how many tests to run\n\n" );
}
static int GetTestCase( const char *name, Type *outType, Type *inType, SaturationMode *sat, RoundingMode *round )
{
int i;
//Find the return type
for( i = 0; i < kTypeCount; i++ )
if( name == strstr( name, gTypeNames[i] ) )
{
*outType = (Type)i;
name += strlen( gTypeNames[i] );
break;
}
if( i == kTypeCount )
return -1;
// Check to see if _sat appears next
*sat = (SaturationMode)0;
for( i = 1; i < kSaturationModeCount; i++ )
if( name == strstr( name, gSaturationNames[i] ) )
{
*sat = (SaturationMode)i;
name += strlen( gSaturationNames[i] );
break;
}
*round = (RoundingMode)0;
for( i = 1; i < kRoundingModeCount; i++ )
if( name == strstr( name, gRoundingModeNames[i] ) )
{
*round = (RoundingMode)i;
name += strlen( gRoundingModeNames[i] );
break;
}
if( *name != '_' )
return -2;
name++;
for( i = 0; i < kTypeCount; i++ )
if( name == strstr( name, gTypeNames[i] ) )
{
*inType = (Type)i;
name += strlen( gTypeNames[i] );
break;
}
if( i == kTypeCount )
return -3;
if( *name != '\0' )
return -4;
return 0;
}
#pragma mark -
#pragma mark OpenCL
test_status InitCL( cl_device_id device )
{
int error, i;
size_t configSize = sizeof( gComputeDevices );
if( (error = clGetDeviceInfo( device, CL_DEVICE_MAX_COMPUTE_UNITS, configSize, &gComputeDevices, NULL )) )
gComputeDevices = 1;
configSize = sizeof( gDeviceFrequency );
if( (error = clGetDeviceInfo( device, CL_DEVICE_MAX_CLOCK_FREQUENCY, configSize, &gDeviceFrequency, NULL )) )
gDeviceFrequency = 0;
cl_device_fp_config floatCapabilities = 0;
if( (error = clGetDeviceInfo(device, CL_DEVICE_SINGLE_FP_CONFIG, sizeof(floatCapabilities), &floatCapabilities, NULL)))
floatCapabilities = 0;
if(0 == (CL_FP_DENORM & floatCapabilities) )
gForceFTZ ^= 1;
if( 0 == (floatCapabilities & CL_FP_ROUND_TO_NEAREST ) )
{
char profileStr[128] = "";
// Verify that we are an embedded profile device
if( (error = clGetDeviceInfo( device, CL_DEVICE_PROFILE, sizeof( profileStr ), profileStr, NULL ) ) )
{
vlog_error( "FAILURE: Could not get device profile: error %d\n", error );
return TEST_FAIL;
}
if( strcmp( profileStr, "EMBEDDED_PROFILE" ) )
{
vlog_error( "FAILURE: non-embedded profile device does not support CL_FP_ROUND_TO_NEAREST\n" );
return TEST_FAIL;
}
if( 0 == (floatCapabilities & CL_FP_ROUND_TO_ZERO ) )
{
vlog_error( "FAILURE: embedded profile device supports neither CL_FP_ROUND_TO_NEAREST or CL_FP_ROUND_TO_ZERO\n" );
return TEST_FAIL;
}
gIsRTZ = 1;
}
else if(is_extension_available(device, "cl_khr_fp64"))
{
gHasDouble = 1;
}
gTestDouble &= gHasDouble;
//detect whether profile of the device is embedded
char profile[1024] = "";
if( (error = clGetDeviceInfo( device, CL_DEVICE_PROFILE, sizeof(profile), profile, NULL ) ) ){}
else if( strstr(profile, "EMBEDDED_PROFILE" ) )
{
gIsEmbedded = 1;
if( !is_extension_available(device, "cles_khr_int64" ) )
gHasLong = 0;
}
gContext = clCreateContext( NULL, 1, &device, notify_callback, NULL, &error );
if( NULL == gContext || error )
{
vlog_error( "clCreateContext failed. (%d)\n", error );
return TEST_FAIL;
}
gQueue = clCreateCommandQueue(gContext, device, 0, &error);
if( NULL == gQueue || error )
{
vlog_error( "clCreateCommandQueue failed. (%d)\n", error );
return TEST_FAIL;
}
//Allocate buffers
//FIXME: use clProtectedArray for guarded allocations?
gIn = malloc( BUFFER_SIZE + 2 * kPageSize );
gAllowZ = malloc( BUFFER_SIZE + 2 * kPageSize );
gRef = malloc( BUFFER_SIZE + 2 * kPageSize );
for( i = 0; i < kCallStyleCount; i++ )
{
gOut[i] = malloc( BUFFER_SIZE + 2 * kPageSize );
if( NULL == gOut[i] )
return TEST_FAIL;
}
// setup input buffers
gInBuffer = clCreateBuffer(gContext, CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR, BUFFER_SIZE, NULL, &error);
if( gInBuffer == NULL || error)
{
vlog_error( "clCreateBuffer failed for input (%d)\n", error );
return TEST_FAIL;
}
// setup output buffers
for( i = 0; i < kCallStyleCount; i++ )
{
gOutBuffers[i] = clCreateBuffer( gContext, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, BUFFER_SIZE, NULL, &error );
if( gOutBuffers[i] == NULL || error )
{
vlog_error( "clCreateArray failed for output (%d)\n", error );
return TEST_FAIL;
}
}
gMTdata = init_genrand( gRandomSeed );
char c[1024];
static const char *no_yes[] = { "NO", "YES" };
vlog( "\nCompute Device info:\n" );
clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(c), c, NULL);
vlog( "\tDevice Name: %s\n", c );
clGetDeviceInfo(device, CL_DEVICE_VENDOR, sizeof(c), c, NULL);
vlog( "\tVendor: %s\n", c );
clGetDeviceInfo(device, CL_DEVICE_VERSION, sizeof(c), c, NULL);
vlog( "\tDevice Version: %s\n", c );
clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_VERSION, sizeof(c), &c, NULL);
vlog( "\tCL C Version: %s\n", c );
clGetDeviceInfo(device, CL_DRIVER_VERSION, sizeof(c), c, NULL);
vlog( "\tDriver Version: %s\n", c );
vlog( "\tProcessing with %ld devices\n", gComputeDevices );
vlog( "\tDevice Frequency: %d MHz\n", gDeviceFrequency );
vlog( "\tSubnormal values supported for floats? %s\n", no_yes[0 != (CL_FP_DENORM & floatCapabilities)] );
vlog( "\tTesting with FTZ mode ON for floats? %s\n", no_yes[0 != gForceFTZ] );
vlog( "\tTesting with default RTZ mode for floats? %s\n", no_yes[0 != gIsRTZ] );
vlog( "\tHas Double? %s\n", no_yes[0 != gHasDouble] );
if( gHasDouble )
vlog( "\tTest Double? %s\n", no_yes[0 != gTestDouble] );
vlog( "\tHas Long? %s\n", no_yes[0 != gHasLong] );
vlog( "\tTesting vector sizes: " );
for( i = gMinVectorSize; i < gMaxVectorSize; i++ )
vlog("\t%d", vectorSizes[i]);
vlog( "\n" );
return TEST_PASS;
}
static int RunKernel( cl_kernel kernel, void *inBuf, void *outBuf, size_t blockCount )
{
// The global dimensions are just the blockCount to execute since we haven't set up multiple queues for multiple devices.
int error;
error = clSetKernelArg(kernel, 0, sizeof( inBuf ), &inBuf);
error |= clSetKernelArg(kernel, 1, sizeof(outBuf), &outBuf);
if( error )
{
vlog_error( "FAILED -- could not set kernel args (%d)\n", error );
return error;
}
if( (error = clEnqueueNDRangeKernel(gQueue, kernel, 1, NULL, &blockCount, NULL, 0, NULL, NULL)))
{
vlog_error( "FAILED -- could not execute kernel (%d)\n", error );
return error;
}
return 0;
}
#if ! defined( __APPLE__ )
void memset_pattern4(void *dest, const void *src_pattern, size_t bytes );
#endif
#if defined( __APPLE__ )
#include <mach/mach_time.h>
#endif
uint64_t GetTime( void );
uint64_t GetTime( void )
{
#if defined( __APPLE__ )
return mach_absolute_time();
#elif defined(_MSC_VER)
return ReadTime();
#else
//mach_absolute_time is a high precision timer with precision < 1 microsecond.
#warning need accurate clock here. Times are invalid.
return 0;
#endif
}
#if defined (_MSC_VER)
/* function is defined in "compat.h" */
#else
double SubtractTime( uint64_t endTime, uint64_t startTime );
double SubtractTime( uint64_t endTime, uint64_t startTime )
{
uint64_t diff = endTime - startTime;
static double conversion = 0.0;
if( 0.0 == conversion )
{
#if defined( __APPLE__ )
mach_timebase_info_data_t info = {0,0};
kern_return_t err = mach_timebase_info( &info );
if( 0 == err )
conversion = 1e-9 * (double) info.numer / (double) info.denom;
#else
// This function consumes output from GetTime() above, and converts the time to secionds.
#warning need accurate ticks to seconds conversion factor here. Times are invalid.
#endif
}
// strictly speaking we should also be subtracting out timer latency here
return conversion * (double) diff;
}
#endif
typedef struct CalcReferenceValuesInfo
{
struct WriteInputBufferInfo *parent; // pointer back to the parent WriteInputBufferInfo struct
cl_kernel kernel; // the kernel for this vector size
cl_program program; // the program for this vector size
cl_uint vectorSize; // the vector size for this callback chain
void *p; // the pointer to mapped result data for this vector size
cl_int result;
}CalcReferenceValuesInfo;
typedef struct WriteInputBufferInfo
{
volatile cl_event calcReferenceValues; // user event which signals when main thread is done calculating reference values
volatile cl_event doneBarrier; // user event which signals when worker threads are done
cl_uint count; // the number of elements in the array
Type outType; // the data type of the conversion result
Type inType; // the data type of the conversion input
volatile int barrierCount;
CalcReferenceValuesInfo calcInfo[kCallStyleCount];
}WriteInputBufferInfo;
cl_uint RoundUpToNextPowerOfTwo( cl_uint x );
cl_uint RoundUpToNextPowerOfTwo( cl_uint x )
{
if( 0 == (x & (x-1)))
return x;
while( x & (x-1) )
x &= x-1;
return x + x;
}
void CL_CALLBACK WriteInputBufferComplete( cl_event, cl_int, void * );
typedef struct DataInitInfo
{
cl_ulong start;
cl_uint size;
Type outType;
Type inType;
SaturationMode sat;
RoundingMode round;
MTdata *d;
}DataInitInfo;
cl_int InitData( cl_uint job_id, cl_uint thread_id, void *p );
cl_int InitData( cl_uint job_id, cl_uint thread_id, void *p )
{
DataInitInfo *info = (DataInitInfo*) p;
gInitFunctions[ info->inType ]( (char*)gIn + job_id * info->size * gTypeSizes[info->inType], info->sat, info->round,
info->outType, info->start + job_id * info->size, info->size, info->d[thread_id] );
return CL_SUCCESS;
}
static void setAllowZ(uint8_t *allow, uint32_t *x, cl_uint count)
{
cl_uint i;
for (i = 0; i < count; ++i)
allow[i] |= (uint8_t)((x[i] & 0x7f800000U) == 0);
}
cl_int PrepareReference( cl_uint job_id, cl_uint thread_id, void *p );
cl_int PrepareReference( cl_uint job_id, cl_uint thread_id, void *p )
{
DataInitInfo *info = (DataInitInfo*) p;
cl_uint count = info->size;
Type inType = info->inType;
Type outType = info->outType;
RoundingMode round = info->round;
size_t j;
Force64BitFPUPrecision();
void *s = (cl_uchar*) gIn + job_id * count * gTypeSizes[info->inType];
void *a = (cl_uchar*) gAllowZ + job_id * count;
void *d = (cl_uchar*) gRef + job_id * count * gTypeSizes[info->outType];
if (outType != inType)
{
//create the reference while we wait
Convert f = gConversions[ outType ][ inType ];
if( info->sat )
f = gSaturatedConversions[ outType ][ inType ];
#if (defined(__arm__) || defined(__aarch64__)) && defined(__GNUC__)
/* ARM VFP doesn't have hardware instruction for converting from 64-bit
* integer to float types, hence GCC ARM uses the floating-point
* emulation code despite which -mfloat-abi setting it is. But the
* emulation code in libgcc.a has only one rounding mode (round to
* nearest even in this case) and ignores the user rounding mode setting
* in hardware. As a result setting rounding modes in hardware won't
* give correct rounding results for type covert from 64-bit integer to
* float using GCC for ARM compiler so for testing different rounding
* modes, we need to use alternative reference function. ARM64 does have
* an instruction, however we cannot guarantee the compiler will use it.
* On all ARM architechures use emulation to calculate reference.*/
switch (round)
{
/* conversions to floating-point type use the current rounding mode.
* The only default floating-point rounding mode supported is round to nearest even
* i.e the current rounding mode will be _rte for floating-point types. */
case kDefaultRoundingMode:
qcom_rm = qcomRTE;
break;
case kRoundToNearestEven:
qcom_rm = qcomRTE;
break;
case kRoundUp:
qcom_rm = qcomRTP;
break;
case kRoundDown:
qcom_rm = qcomRTN;
break;
case kRoundTowardZero:
qcom_rm = qcomRTZ;
break;
default:
vlog_error("ERROR: undefined rounding mode %d\n", round);
break;
}
qcom_sat = info->sat;
#endif
RoundingMode oldRound = set_round( round, outType );
f( d, s, count );
set_round( oldRound, outType );
// Decide if we allow a zero result in addition to the correctly rounded one
memset(a, 0, count);
if (gForceFTZ) {
if (inType == kfloat)
setAllowZ((uint8_t*)a, (uint32_t*)s, count);
if (outType == kfloat)
setAllowZ((uint8_t*)a, (uint32_t*)d, count);
}
}
else
{
// Copy the input to the reference
memcpy(d, s, info->size * gTypeSizes[inType]);
}
//Patch up NaNs conversions to integer to zero -- these can be converted to any integer
if( info->outType != kfloat && info->outType != kdouble )
{
if( inType == kfloat )
{
float *inp = (float*) s;
for( j = 0; j < count; j++ )
{
if( isnan( inp[j] ) )
memset( (char*) d + j * gTypeSizes[ outType ], 0, gTypeSizes[ outType ] );
}
}
if( inType == kdouble )
{
double *inp = (double*) s;
for( j = 0; j < count; j++ )
{
if( isnan( inp[j] ) )
memset( (char*) d + j * gTypeSizes[ outType ], 0, gTypeSizes[ outType ] );
}
}
}
else if( inType == kfloat || inType == kdouble )
{ // outtype and intype is float or double. NaN conversions for float <-> double can be any NaN
if( inType == kfloat && outType == kdouble )
{
float *inp = (float*) s;
for( j = 0; j < count; j++ )
{
if( isnan( inp[j] ) )
((double*) d)[j] = NAN;
}
}
if( inType == kdouble && outType == kfloat )
{
double *inp = (double*) s;
for( j = 0; j < count; j++ )
{
if( isnan( inp[j] ) )
((float*) d)[j] = NAN;
}
}
}
return CL_SUCCESS;
}
static int DoTest( cl_device_id device, Type outType, Type inType, SaturationMode sat, RoundingMode round, MTdata d )
{
#ifdef __APPLE__
cl_ulong wall_start = mach_absolute_time();
#endif
DataInitInfo init_info = { 0, 0, outType, inType, sat, round, NULL };
WriteInputBufferInfo writeInputBufferInfo;
int vectorSize;
int error = 0;
cl_uint threads = GetThreadCount();
uint64_t i;
gTestCount++;
size_t blockCount = BUFFER_SIZE / MAX( gTypeSizes[ inType ], gTypeSizes[ outType ] );
size_t step = blockCount;
uint64_t lastCase = 1ULL << (8*gTypeSizes[ inType ]);
cl_event writeInputBuffer = NULL;
memset( &writeInputBufferInfo, 0, sizeof( writeInputBufferInfo ) );
init_info.d = (MTdata*)malloc( threads * sizeof( MTdata ) );
if( NULL == init_info.d )
{
vlog_error( "ERROR: Unable to allocate storage for random number generator!\n" );
return -1;
}
for( i = 0; i < threads; i++ )
{
init_info.d[i] = init_genrand( genrand_int32( d ) );
if( NULL == init_info.d[i] )
{
vlog_error( "ERROR: Unable to allocate storage for random number generator!\n" );
return -1;
}
}
writeInputBufferInfo.outType = outType;
writeInputBufferInfo.inType = inType;
for( vectorSize = gMinVectorSize; vectorSize < gMaxVectorSize; vectorSize++)
{
writeInputBufferInfo.calcInfo[vectorSize].program = MakeProgram( outType, inType, sat, round, vectorSize,
&writeInputBufferInfo.calcInfo[vectorSize].kernel );
if( NULL == writeInputBufferInfo.calcInfo[vectorSize].program )
{
gFailCount++;
return -1;
}
if( NULL == writeInputBufferInfo.calcInfo[vectorSize].kernel )
{
gFailCount++;
vlog_error( "\t\tFAILED -- Failed to create kernel.\n" );
return -2;
}
writeInputBufferInfo.calcInfo[vectorSize].parent = &writeInputBufferInfo;
writeInputBufferInfo.calcInfo[vectorSize].vectorSize = vectorSize;
writeInputBufferInfo.calcInfo[vectorSize].result = -1;
}
if( gSkipTesting )
goto exit;
// Patch up rounding mode if default is RTZ
// We leave the part above in default rounding mode so that the right kernel is compiled.
if( round == kDefaultRoundingMode && gIsRTZ && (outType == kfloat) )
init_info.round = round = kRoundTowardZero;
// Figure out how many elements are in a work block
// we handle 64-bit types a bit differently.
if( 8*gTypeSizes[ inType ] > 32 )
lastCase = 0x100000000ULL;
if ( !gWimpyMode && gIsEmbedded )
step = blockCount * EMBEDDED_REDUCTION_FACTOR;
if ( gWimpyMode )
step = (size_t)blockCount * (size_t)gWimpyReductionFactor;
vlog( "Testing... " );
fflush(stdout);
for( i = 0; i < (uint64_t)lastCase; i += step )
{
if( 0 == ( i & ((lastCase >> 3) -1))) {
vlog(".");
fflush(stdout);
}
cl_uint count = (uint32_t) MIN( blockCount, lastCase - i );
writeInputBufferInfo.count = count;
// Crate a user event to represent the status of the reference value computation completion
writeInputBufferInfo.calcReferenceValues = clCreateUserEvent( gContext, &error);
if( error || NULL == writeInputBufferInfo.calcReferenceValues )
{
vlog_error( "ERROR: Unable to create user event. (%d)\n", error );
gFailCount++;
goto exit;
}
// retain for consumption by MapOutputBufferComplete
for( vectorSize = gMinVectorSize; vectorSize < gMaxVectorSize; vectorSize++)
{
if( (error = clRetainEvent(writeInputBufferInfo.calcReferenceValues) ))
{
vlog_error( "ERROR: Unable to retain user event. (%d)\n", error );
gFailCount++;
goto exit;
}
}
// Crate a user event to represent when the callbacks are done verifying correctness
writeInputBufferInfo.doneBarrier = clCreateUserEvent( gContext, &error);
if( error || NULL == writeInputBufferInfo.calcReferenceValues )
{
vlog_error( "ERROR: Unable to create user event for barrier. (%d)\n", error );
gFailCount++;
goto exit;
}
// retain for use by the callback that calls this
if( (error = clRetainEvent(writeInputBufferInfo.doneBarrier) ))
{
vlog_error( "ERROR: Unable to retain user event doneBarrier. (%d)\n", error );
gFailCount++;
goto exit;
}
// Call this in a multithreaded manner
// gInitFunctions[ inType ]( gIn, sat, round, outType, i, count, d );
cl_uint chunks = RoundUpToNextPowerOfTwo(threads) * 2;
init_info.start = i;
init_info.size = count / chunks;
if( init_info.size < 16384 )
{
chunks = RoundUpToNextPowerOfTwo(threads);
init_info.size = count / chunks;
if( init_info.size < 16384 )
{
init_info.size = count;
chunks = 1;
}
}
ThreadPool_Do(InitData, chunks, &init_info);
// Copy the results to the device
writeInputBuffer = NULL;
if( (error = clEnqueueWriteBuffer(gQueue, gInBuffer, CL_FALSE, 0, count * gTypeSizes[inType], gIn, 0, NULL, &writeInputBuffer )))
{
vlog_error( "ERROR: clEnqueueWriteBuffer failed. (%d)\n", error );
gFailCount++;
goto exit;
}
// Setup completion callback for the write, which will enqueue the rest of the work
// This is somewhat gratuitous. Because this is an in order queue, we didn't really need to
// do this work in a callback. We could have done it from the main thread. Here we are
// verifying that the implementation can enqueue work from a callback, while at the same time
// also checking to make sure that the conversions work.
//
// Because the verification code is also moved to a callback, it is hoped that implementations will
// achieve a test performance improvement because they can verify the results in parallel. If the
// implementation serializes callbacks however, that won't happen. Consider it some motivation
// to do the right thing! :-)
if( (error = clSetEventCallback( writeInputBuffer, CL_COMPLETE, WriteInputBufferComplete, &writeInputBufferInfo)) )
{
vlog_error( "ERROR: clSetEventCallback failed. (%d)\n", error );
gFailCount++;
goto exit;
}
// The event can't be destroyed until the callback is called, so we can release it now.
if( (error = clReleaseEvent(writeInputBuffer) ))
{
vlog_error( "ERROR: clReleaseEvent failed. (%d)\n", error );
gFailCount++;
goto exit;
}
// Make sure the work is actually running, so we don't deadlock
if( (error = clFlush( gQueue ) ) )
{
vlog_error( "clFlush failed with error %d\n", error );
gFailCount++;
goto exit;
}
ThreadPool_Do(PrepareReference, chunks, &init_info);
// signal we are done calculating the reference results
if( (error = clSetUserEventStatus( writeInputBufferInfo.calcReferenceValues, CL_COMPLETE ) ) )
{
vlog_error( "Error: Failed to set user event status to CL_COMPLETE: %d\n", error );
gFailCount++;
goto exit;
}
// Wait for the event callbacks to finish verifying correctness.
if( (error = clWaitForEvents( 1, (cl_event*) &writeInputBufferInfo.doneBarrier ) ))
{
vlog_error( "Error: Failed to wait for barrier: %d\n", error );
gFailCount++;
goto exit;
}
if( (error = clReleaseEvent(writeInputBufferInfo.calcReferenceValues ) ))
{
vlog_error( "Error: Failed to release calcReferenceValues: %d\n", error );
gFailCount++;
goto exit;
}
if( (error = clReleaseEvent(writeInputBufferInfo.doneBarrier ) ))
{
vlog_error( "Error: Failed to release done barrier: %d\n", error );
gFailCount++;
goto exit;
}
for( vectorSize = gMinVectorSize; vectorSize < gMaxVectorSize; vectorSize++)
{
if( ( error = writeInputBufferInfo.calcInfo[ vectorSize ].result ))
{
switch( inType )
{
case kuchar:
case kchar:
vlog( "Input value: 0x%2.2x ", ((unsigned char*)gIn)[error - 1] );
break;
case kushort:
case kshort:
vlog( "Input value: 0x%4.4x ", ((unsigned short*)gIn)[error - 1] );
break;
case kuint:
case kint:
vlog( "Input value: 0x%8.8x ", ((unsigned int*)gIn)[error - 1] );
break;
case kfloat:
vlog( "Input value: %a ", ((float*)gIn)[error - 1] );
break;
break;
case kulong:
case klong:
vlog( "Input value: 0x%16.16llx ", ((unsigned long long*)gIn)[error - 1] );
break;
case kdouble:
vlog( "Input value: %a ", ((double*)gIn)[error - 1]);
break;
default:
vlog_error( "Internal error at %s: %d\n", __FILE__, __LINE__ );
abort();
break;
}
// tell the user which conversion it was.
if( 0 == vectorSize )
vlog( " (implicit scalar conversion from %s to %s)\n", gTypeNames[ inType ], gTypeNames[ outType ] );
else
vlog( " (convert_%s%s%s%s( %s%s ))\n", gTypeNames[outType], sizeNames[vectorSize], gSaturationNames[ sat ],
gRoundingModeNames[ round ], gTypeNames[inType], sizeNames[vectorSize] );
gFailCount++;
goto exit;
}
}
}
log_info( "done.\n" );
if( gTimeResults )
{
//Kick off tests for the various vector lengths
for( vectorSize = gMinVectorSize; vectorSize < gMaxVectorSize; vectorSize++)
{
size_t workItemCount = blockCount / vectorSizes[vectorSize];
if( vectorSizes[vectorSize] * gTypeSizes[outType] < 4 )
workItemCount /= 4 / (vectorSizes[vectorSize] * gTypeSizes[outType]);
double sum = 0.0;
double bestTime = INFINITY;
cl_uint k;
for( k = 0; k < PERF_LOOP_COUNT; k++ )
{
uint64_t startTime = GetTime();
if( (error = RunKernel( writeInputBufferInfo.calcInfo[vectorSize].kernel, gInBuffer, gOutBuffers[ vectorSize ], workItemCount )) )
{
gFailCount++;
goto exit;
}
// Make sure OpenCL is done
if( (error = clFinish(gQueue) ) )
{
vlog_error( "Error %d at clFinish\n", error );
goto exit;
}
uint64_t endTime = GetTime();
double time = SubtractTime( endTime, startTime );
sum += time;
if( time < bestTime )
bestTime = time;
}
if( gReportAverageTimes )
bestTime = sum / PERF_LOOP_COUNT;
double clocksPerOp = bestTime * (double) gDeviceFrequency * gComputeDevices * gSimdSize * 1e6 / (workItemCount * vectorSizes[vectorSize]);
if( 0 == vectorSize )
vlog_perf( clocksPerOp, LOWER_IS_BETTER, "clocks / element", "implicit convert %s -> %s", gTypeNames[ inType ], gTypeNames[ outType ] );
else
vlog_perf( clocksPerOp, LOWER_IS_BETTER, "clocks / element", "convert_%s%s%s%s( %s%s )", gTypeNames[ outType ], sizeNames[vectorSize], gSaturationNames[ sat ], gRoundingModeNames[round], gTypeNames[inType], sizeNames[vectorSize] );
}
}
if( gWimpyMode )
vlog( "\tWimp pass" );
else
vlog( "\tpassed" );
#ifdef __APPLE__
// record the run time
vlog( "\t(%f s)", 1e-9 * ( mach_absolute_time() - wall_start ) );
#endif
vlog( "\n\n" );
fflush( stdout );
exit:
//clean up
for( vectorSize = gMinVectorSize; vectorSize < gMaxVectorSize; vectorSize++)
{
clReleaseProgram( writeInputBufferInfo.calcInfo[vectorSize].program );
clReleaseKernel( writeInputBufferInfo.calcInfo[vectorSize].kernel );
}
if( init_info.d )
{
for( i = 0; i < threads; i++ )
free_mtdata(init_info.d[i]);
free(init_info.d);
}
return error;
}
void CL_CALLBACK MapResultValuesComplete( cl_event e, cl_int status, void *data );
// Note: not called reentrantly
void CL_CALLBACK WriteInputBufferComplete( cl_event e, cl_int status, void *data )
{
WriteInputBufferInfo *info = (WriteInputBufferInfo*) data;
cl_uint count = info->count;
int vectorSize;
if( CL_SUCCESS != status )
{
vlog_error( "ERROR: WriteInputBufferComplete calback failed with status: %d\n", status );
gFailCount++;
return;
}
info->barrierCount = gMaxVectorSize - gMinVectorSize;
// now that we know that the write buffer is complete, enqueue callbacks to wait for the main thread to
// finish calculating the reference results.
for( vectorSize = gMinVectorSize; vectorSize < gMaxVectorSize; vectorSize++)
{
size_t workItemCount = (count + vectorSizes[vectorSize] - 1) / ( vectorSizes[vectorSize]);
cl_event mapComplete = NULL;
if( (status = RunKernel( info->calcInfo[ vectorSize ].kernel, gInBuffer, gOutBuffers[ vectorSize ], workItemCount )) )
{
gFailCount++;
return;
}
info->calcInfo[vectorSize].p = clEnqueueMapBuffer( gQueue, gOutBuffers[ vectorSize ], CL_FALSE, CL_MAP_READ | CL_MAP_WRITE,
0, count * gTypeSizes[ info->outType ], 0, NULL, &mapComplete, &status);
{
if( status )
{
vlog_error( "ERROR: WriteInputBufferComplete calback failed with status: %d\n", status );
gFailCount++;
return;
}
}
if( (status = clSetEventCallback( mapComplete, CL_COMPLETE, MapResultValuesComplete, info->calcInfo + vectorSize)))
{
vlog_error( "ERROR: WriteInputBufferComplete calback failed with status: %d\n", status );
gFailCount++;
return;
}
if( (status = clReleaseEvent(mapComplete)))
{
vlog_error( "ERROR: clReleaseEvent calback failed in WriteInputBufferComplete for vector size %d with status: %d\n", vectorSize, status );
gFailCount++;
return;
}
}
// Make sure the work starts moving -- otherwise we may deadlock
if( (status = clFlush(gQueue)))
{
vlog_error( "ERROR: WriteInputBufferComplete calback failed with status: %d\n", status );
gFailCount++;
return;
}
// e was already released by the main thread. It should be destroyed automatically soon after we exit.
}
void CL_CALLBACK CalcReferenceValuesComplete( cl_event e, cl_int status, void *data );
// Note: May be called reentrantly
void CL_CALLBACK MapResultValuesComplete( cl_event e, cl_int status, void *data )
{
CalcReferenceValuesInfo *info = (CalcReferenceValuesInfo*) data;
cl_event calcReferenceValues = info->parent->calcReferenceValues;
if( CL_SUCCESS != status )
{
vlog_error( "ERROR: MapResultValuesComplete calback failed with status: %d\n", status );
gFailCount++; // not thread safe -- being lazy here
clReleaseEvent(calcReferenceValues);
return;
}
// we know that the map is done, wait for the main thread to finish calculating the reference values
if( (status = clSetEventCallback( calcReferenceValues, CL_COMPLETE, CalcReferenceValuesComplete, data )))
{
vlog_error( "ERROR: clSetEventCallback failed in MapResultValuesComplete with status: %d\n", status );
gFailCount++; // not thread safe -- being lazy here
}
// this thread no longer needs its reference to info->calcReferenceValues, so release it
if( (status = clReleaseEvent(calcReferenceValues) ))
{
vlog_error( "ERROR: clReleaseEvent(info->calcReferenceValues) failed with status: %d\n", status );
gFailCount++; // not thread safe -- being lazy here
}
// no need to flush since we didn't enqueue anything
// e was already released by WriteInputBufferComplete. It should be destroyed automatically soon after we exit.
}
void CL_CALLBACK CalcReferenceValuesComplete( cl_event e, cl_int status, void *data )
{
CalcReferenceValuesInfo *info = (CalcReferenceValuesInfo*) data;
cl_uint vectorSize = info->vectorSize;
cl_uint count = info->parent->count;
Type outType = info->parent->outType; // the data type of the conversion result
Type inType = info->parent->inType; // the data type of the conversion input
size_t j;
cl_int error;
cl_event doneBarrier = info->parent->doneBarrier;
// report spurious error condition
if( CL_SUCCESS != status )
{
vlog_error( "ERROR: CalcReferenceValuesComplete did not succeed! (%d)\n", status );
gFailCount++; // lazy about thread safety here
return;
}
// Now we know that both results have been mapped back from the device, and the
// main thread is done calculating the reference results. It is now time to check
// the results.
// verify results
void *mapped = info->p;
//Patch up NaNs conversions to integer to zero -- these can be converted to any integer
if( outType != kfloat && outType != kdouble )
{
if( inType == kfloat )
{
float *inp = (float*) gIn;
for( j = 0; j < count; j++ )
{
if( isnan( inp[j] ) )
memset( (char*) mapped + j * gTypeSizes[ outType ], 0, gTypeSizes[ outType ] );
}
}
if( inType == kdouble )
{
double *inp = (double*) gIn;
for( j = 0; j < count; j++ )
{
if( isnan( inp[j] ) )
memset( (char*) mapped + j * gTypeSizes[ outType ], 0, gTypeSizes[ outType ] );
}
}
}
else if( inType == kfloat || inType == kdouble )
{ // outtype and intype is float or double. NaN conversions for float <-> double can be any NaN
if( inType == kfloat && outType == kdouble )
{
float *inp = (float*) gIn;
double *outp = (double*) mapped;
for( j = 0; j < count; j++ )
{
if( isnan( inp[j] ) && isnan(outp[j]) )
outp[j] = NAN;
}
}
if( inType == kdouble && outType == kfloat )
{
double *inp = (double*) gIn;
float *outp = (float*) mapped;
for( j = 0; j < count; j++ )
{
if( isnan( inp[j] ) && isnan(outp[j]) )
outp[j] = NAN;
}
}
}
if( memcmp( mapped, gRef, count * gTypeSizes[ outType ] ) )
info->result = gCheckResults[outType]( mapped, gRef, gAllowZ, count, vectorSizes[vectorSize] );
else
info->result = 0;
// Fill the output buffer with junk and release it
{
cl_uint pattern = 0xffffdead;
memset_pattern4(mapped, &pattern, count * gTypeSizes[outType]);
if((error = clEnqueueUnmapMemObject(gQueue, gOutBuffers[ vectorSize ], mapped, 0, NULL, NULL)))
{
vlog_error( "ERROR: clEnqueueUnmapMemObject failed in CalcReferenceValuesComplete (%d)\n", error );
gFailCount++;
}
}
if( 1 == ThreadPool_AtomicAdd( &info->parent->barrierCount, -1) )
{
if( (status = clSetUserEventStatus( doneBarrier, CL_COMPLETE) ))
{
vlog_error( "ERROR: clSetUserEventStatus failed in CalcReferenceValuesComplete (err: %d). We're probably going to deadlock.\n", status );
gFailCount++;
return;
}
if( (status = clReleaseEvent( doneBarrier ) ) )
{
vlog_error( "ERROR: clReleaseEvent failed in CalcReferenceValuesComplete (err: %d).\n", status );
gFailCount++;
return;
}
}
// e was already released by WriteInputBufferComplete. It should be destroyed automatically soon after
// all the calls to CalcReferenceValuesComplete exit.
}
static cl_program MakeProgram( Type outType, Type inType, SaturationMode sat, RoundingMode round, int vectorSize, cl_kernel *outKernel )
{
cl_program program;
char testName[256];
int error = 0;
const char **strings;
size_t stringCount = 0;
// Create the program. This is a bit complicated because we are trying to avoid byte and short stores.
if (0 == vectorSize)
{
char inName[32];
char outName[32];
const char *programSource[] =
{
"", // optional pragma
"__kernel void ", testName, "( __global ", inName, " *src, __global ", outName, " *dest )\n"
"{\n"
" size_t i = get_global_id(0);\n"
" dest[i] = src[i];\n"
"}\n"
};
stringCount = sizeof(programSource) / sizeof(programSource[0]);
strings = programSource;
if (outType == kdouble || inType == kdouble)
programSource[0] = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
//create the type name
strncpy(inName, gTypeNames[inType], sizeof(inName));
strncpy(outName, gTypeNames[outType], sizeof(outName));
sprintf(testName, "test_implicit_%s_%s", outName, inName);
vlog("Building implicit %s -> %s conversion test\n", gTypeNames[inType], gTypeNames[outType]);
fflush(stdout);
}
else
{
int vectorSizetmp = vectorSizes[vectorSize];
char convertString[128];
char inName[32];
char outName[32];
const char *programSource[] =
{
"", // optional pragma
"__kernel void ", testName, "( __global ", inName, " *src, __global ", outName, " *dest )\n"
"{\n"
" size_t i = get_global_id(0);\n"
" dest[i] = ", convertString, "( src[i] );\n"
"}\n"
};
const char *programSourceV3[] =
{
"", // optional pragma
"__kernel void ", testName, "( __global ", inName, " *src, __global ", outName, " *dest )\n"
"{\n"
" size_t i = get_global_id(0);\n"
" if( i + 1 < get_global_size(0))\n"
" vstore3( ", convertString, "( vload3( i, src)), i, dest );\n"
" else\n"
" {\n"
" ", inName, "3 in;\n"
" ", outName, "3 out;\n"
" if( 0 == (i & 1) )\n"
" in.y = src[3*i+1];\n"
" in.x = src[3*i];\n"
" out = ", convertString, "( in ); \n"
" dest[3*i] = out.x;\n"
" if( 0 == (i & 1) )\n"
" dest[3*i+1] = out.y;\n"
" }\n"
"}\n"
};
stringCount = 3 == vectorSizetmp ? sizeof(programSourceV3) / sizeof(programSourceV3[0]) :
sizeof(programSource) / sizeof(programSource[0]);
strings = 3 == vectorSizetmp ? programSourceV3 : programSource;
if (outType == kdouble || inType == kdouble) {
programSource[0] = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
programSourceV3[0] = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
}
//create the type name
switch (vectorSizetmp)
{
case 1:
strncpy(inName, gTypeNames[inType], sizeof(inName));
strncpy(outName, gTypeNames[outType], sizeof(outName));
snprintf(convertString, sizeof(convertString), "convert_%s%s%s", outName, gSaturationNames[sat], gRoundingModeNames[round]);
snprintf(testName, 256, "test_%s_%s", convertString, inName);
vlog("Building %s( %s ) test\n", convertString, inName);
break;
case 3:
strncpy(inName, gTypeNames[inType], sizeof(inName));
strncpy(outName, gTypeNames[outType], sizeof(outName));
snprintf(convertString, sizeof(convertString), "convert_%s3%s%s", outName, gSaturationNames[sat], gRoundingModeNames[round]);
snprintf(testName, 256, "test_%s_%s3", convertString, inName);
vlog("Building %s( %s3 ) test\n", convertString, inName);
break;
default:
snprintf(inName, sizeof(inName), "%s%d", gTypeNames[inType], vectorSizetmp);
snprintf(outName, sizeof(outName), "%s%d", gTypeNames[outType], vectorSizetmp);
snprintf(convertString, sizeof(convertString), "convert_%s%s%s", outName, gSaturationNames[sat], gRoundingModeNames[round]);
snprintf(testName, 256, "test_%s_%s", convertString, inName);
vlog("Building %s( %s ) test\n", convertString, inName);
break;
}
fflush(stdout);
}
*outKernel = NULL;
const char *flags = NULL;
if( gForceFTZ )
flags = "-cl-denorms-are-zero";
// build it
error = create_single_kernel_helper(gContext, &program, outKernel, (cl_uint)stringCount, strings, testName, flags);
if (error)
{
char buffer[2048] = "";
vlog_error("Failed to build kernel/program.\n", error);
clReleaseProgram(program);
return NULL;
}
return program;
}