You can not select more than 25 topics
Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
1683 lines
60 KiB
1683 lines
60 KiB
//
|
|
// 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;
|
|
}
|