| // |
| // 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 <string.h> |
| #include <stdio.h> |
| |
| #if !defined(_WIN32) |
| #include <libgen.h> |
| #include <sys/param.h> |
| #endif |
| |
| #include "mingw_compat.h" |
| #if defined (__MINGW32__) |
| #include <sys/param.h> |
| #endif |
| |
| #include <time.h> |
| #include "errorHelpers.h" |
| #include "harness/compat.h" |
| #include "harness/mt19937.h" |
| #include "harness/kernelHelpers.h" |
| #include "harness/rounding_mode.h" |
| #include "harness/fpcontrol.h" |
| #include "harness/testHarness.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 (_WIN32) |
| #include <string.h> |
| #endif |
| |
| #if defined(_MSC_VER) && (defined(_M_IX86) || defined(_M_X64)) |
| #include <emmintrin.h> |
| #endif |
| |
| #if defined(__PPC__) |
| // Global varaiable used to hold the FPU control register state. The FPSCR register can not |
| // be used because not all Power implementations retain or observed the NI (non-IEEE |
| // mode) bit. |
| __thread fpu_control_t fpu_control = 0; |
| #endif |
| |
| #ifndef MAXPATHLEN |
| #define MAXPATHLEN 2048 |
| #endif |
| |
| char appName[ MAXPATHLEN ] = ""; |
| cl_context gContext = NULL; |
| cl_command_queue gQueue = NULL; |
| cl_program gProgram[5] = { NULL, NULL, NULL, NULL, NULL }; |
| cl_program gProgram_double[5] = { NULL, NULL, NULL, NULL, NULL }; |
| int gForceFTZ = 0; |
| int gSeed = 0; |
| int gSeedSpecified = 0; |
| int gHasDouble = 0; |
| MTdata gMTdata = NULL; |
| int gSkipNanInf = 0; |
| int gIgnoreZeroSign = 0; |
| |
| cl_mem bufA = NULL; |
| cl_mem bufB = NULL; |
| cl_mem bufC = NULL; |
| cl_mem bufD = NULL; |
| cl_mem bufE = NULL; |
| cl_mem bufC_double = NULL; |
| cl_mem bufD_double = NULL; |
| float *buf1, *buf2, *buf3, *buf4, *buf5, *buf6; |
| float *correct[8]; |
| int *skipTest[8]; |
| |
| double *buf3_double, *buf4_double, *buf5_double, *buf6_double; |
| double *correct_double[8]; |
| |
| static const char **gArgList; |
| static size_t gArgCount; |
| |
| #define BUFFER_SIZE (1024*1024) |
| |
| |
| static int ParseArgs( int argc, const char **argv ); |
| static void PrintUsage( void ); |
| test_status InitCL( cl_device_id device ); |
| static void ReleaseCL( void ); |
| static int RunTest( int testNumber ); |
| static int RunTest_Double( int testNumber ); |
| |
| #if defined(__ANDROID__) |
| #define nanf( X ) strtof( "NAN", ( char ** ) NULL ) |
| #define nan( X ) strtod( "NAN", ( char ** ) NULL ) |
| #endif |
| |
| #if defined(_MSC_VER) && (defined(_M_IX86) || defined(_M_X64)) |
| // defeat x87 on MSVC |
| float sse_add(float x, float y) |
| { |
| volatile float a = x; |
| volatile float b = y; |
| |
| // defeat x87 |
| __m128 va = _mm_set_ss( (float) a ); |
| __m128 vb = _mm_set_ss( (float) b ); |
| va = _mm_add_ss( va, vb ); |
| _mm_store_ss( (float*) &a, va ); |
| return a; |
| } |
| |
| double sse_add_sd(double x, double y) |
| { |
| volatile double a = x; |
| volatile double b = y; |
| |
| // defeat x87 |
| __m128d va = _mm_set_sd( (double) a ); |
| __m128d vb = _mm_set_sd( (double) b ); |
| va = _mm_add_sd( va, vb ); |
| _mm_store_sd( (double*) &a, va ); |
| return a; |
| } |
| |
| float sse_sub(float x, float y) |
| { |
| volatile float a = x; |
| volatile float b = y; |
| |
| // defeat x87 |
| __m128 va = _mm_set_ss( (float) a ); |
| __m128 vb = _mm_set_ss( (float) b ); |
| va = _mm_sub_ss( va, vb ); |
| _mm_store_ss( (float*) &a, va ); |
| return a; |
| } |
| |
| double sse_sub_sd(double x, double y) |
| { |
| volatile double a = x; |
| volatile double b = y; |
| |
| // defeat x87 |
| __m128d va = _mm_set_sd( (double) a ); |
| __m128d vb = _mm_set_sd( (double) b ); |
| va = _mm_sub_sd( va, vb ); |
| _mm_store_sd( (double*) &a, va ); |
| return a; |
| } |
| |
| float sse_mul(float x, float y) |
| { |
| volatile float a = x; |
| volatile float b = y; |
| |
| // defeat x87 |
| __m128 va = _mm_set_ss( (float) a ); |
| __m128 vb = _mm_set_ss( (float) b ); |
| va = _mm_mul_ss( va, vb ); |
| _mm_store_ss( (float*) &a, va ); |
| return a; |
| } |
| |
| double sse_mul_sd(double x, double y) |
| { |
| volatile double a = x; |
| volatile double b = y; |
| |
| // defeat x87 |
| __m128d va = _mm_set_sd( (double) a ); |
| __m128d vb = _mm_set_sd( (double) b ); |
| va = _mm_mul_sd( va, vb ); |
| _mm_store_sd( (double*) &a, va ); |
| return a; |
| } |
| #endif |
| |
| #ifdef __PPC__ |
| float ppc_mul(float a, float b) |
| { |
| float p; |
| |
| if (gForceFTZ) { |
| // Flush input a to zero if it is sub-normal |
| if (fabsf(a) < FLT_MIN) { |
| a = copysignf(0.0, a); |
| } |
| // Flush input b to zero if it is sub-normal |
| if (fabsf(b) < FLT_MIN) { |
| b = copysignf(0.0, b); |
| } |
| // Perform multiply |
| p = a * b; |
| // Flush the product if it is a sub-normal |
| if (fabs((double)a * (double)b) < FLT_MIN) { |
| p = copysignf(0.0, p); |
| } |
| } else { |
| p = a * b; |
| } |
| return p; |
| } |
| #endif |
| |
| int test_contractions_float_0(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest(0); |
| } |
| |
| int test_contractions_float_1(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest(1); |
| } |
| |
| int test_contractions_float_2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest(2); |
| } |
| |
| int test_contractions_float_3(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest(3); |
| } |
| |
| int test_contractions_float_4(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest(4); |
| } |
| |
| int test_contractions_float_5(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest(5); |
| } |
| |
| int test_contractions_float_6(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest(6); |
| } |
| |
| int test_contractions_float_7(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest(7); |
| } |
| |
| int test_contractions_double_0(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest_Double(0); |
| } |
| |
| int test_contractions_double_1(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest_Double(1); |
| } |
| |
| int test_contractions_double_2(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest_Double(2); |
| } |
| |
| int test_contractions_double_3(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest_Double(3); |
| } |
| |
| int test_contractions_double_4(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest_Double(4); |
| } |
| |
| int test_contractions_double_5(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest_Double(5); |
| } |
| |
| int test_contractions_double_6(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest_Double(6); |
| } |
| |
| int test_contractions_double_7(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| return RunTest_Double(7); |
| } |
| |
| test_definition test_list[] = { |
| ADD_TEST( contractions_float_0 ), |
| ADD_TEST( contractions_float_1 ), |
| ADD_TEST( contractions_float_2 ), |
| ADD_TEST( contractions_float_3 ), |
| ADD_TEST( contractions_float_4 ), |
| ADD_TEST( contractions_float_5 ), |
| ADD_TEST( contractions_float_6 ), |
| ADD_TEST( contractions_float_7 ), |
| ADD_TEST( contractions_double_0 ), |
| ADD_TEST( contractions_double_1 ), |
| ADD_TEST( contractions_double_2 ), |
| ADD_TEST( contractions_double_3 ), |
| ADD_TEST( contractions_double_4 ), |
| ADD_TEST( contractions_double_5 ), |
| ADD_TEST( contractions_double_6 ), |
| ADD_TEST( contractions_double_7 ), |
| }; |
| |
| const int test_num = ARRAY_SIZE( test_list ); |
| |
| int main( int argc, const char **argv ) |
| { |
| argc = parseCustomParam(argc, argv); |
| if (argc == -1) |
| { |
| return -1; |
| } |
| |
| int error = ParseArgs( argc, argv ); |
| |
| if( !error ) |
| { |
| error = runTestHarnessWithCheck( gArgCount, gArgList, test_num, test_list, true, 0, InitCL ); |
| } |
| |
| if( gQueue ) |
| { |
| int flush_error = clFinish( gQueue ); |
| if( flush_error ) |
| log_error( "clFinish failed: %d\n", flush_error ); |
| } |
| |
| ReleaseCL(); |
| free( gArgList ); |
| |
| return error; |
| } |
| |
| |
| |
| static int ParseArgs( int argc, const char **argv ) |
| { |
| gArgList = (const char **)calloc( argc, sizeof( char*) ); |
| |
| if( NULL == gArgList ) |
| { |
| vlog_error( "Failed to allocate memory for argList\n" ); |
| return 1; |
| } |
| |
| gArgList[0] = argv[0]; |
| gArgCount = 1; |
| |
| int length_of_seed = 0; |
| |
| { // Extract the app name |
| strncpy( appName, argv[0], MAXPATHLEN ); |
| |
| #if (defined( __APPLE__ ) || defined(__linux__) || defined(__MINGW32__)) |
| char baseName[MAXPATHLEN]; |
| char *base = NULL; |
| strncpy( baseName, argv[0], MAXPATHLEN ); |
| 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 |
| } |
| |
| for( int i = 1; i < argc; i++ ) |
| { |
| const char *arg = argv[i]; |
| if( NULL == arg ) |
| break; |
| |
| if( arg[0] == '-' ) |
| { |
| while( arg[1] != '\0' ) |
| { |
| arg++; |
| switch( *arg ) |
| { |
| case 'h': |
| PrintUsage(); |
| return -1; |
| |
| case 's': |
| arg++; |
| gSeed = atoi( arg ); |
| while (arg[length_of_seed] >='0' && arg[length_of_seed]<='9') |
| length_of_seed++; |
| gSeedSpecified = 1; |
| arg+=length_of_seed-1; |
| break; |
| |
| case 'z': |
| gForceFTZ ^= 1; |
| break; |
| |
| default: |
| vlog( " <-- unknown flag: %c (0x%2.2x)\n)", *arg, *arg ); |
| PrintUsage(); |
| return -1; |
| } |
| } |
| } |
| else |
| { |
| gArgList[gArgCount] = arg; |
| gArgCount++; |
| } |
| } |
| vlog( "\n\nTest binary built %s %s\n", __DATE__, __TIME__ ); |
| |
| PrintArch(); |
| |
| return 0; |
| } |
| |
| static void PrintUsage( void ) |
| { |
| vlog( "%s [-z]: <optional: test names>\n", appName ); |
| vlog( "\tOptions:\n" ); |
| vlog( "\t\t-z\tToggle FTZ mode (Section 6.5.3) for all functions. (Set by device capabilities by default.)\n" ); |
| vlog( "\t\t-sNUMBER set random seed.\n"); |
| vlog( "\n" ); |
| vlog( "\tTest names:\n" ); |
| for( int i = 0; i < test_num; i++ ) |
| { |
| vlog( "\t\t%s\n", test_list[i].name ); |
| } |
| } |
| |
| const char *sizeNames[] = { "float", "float2", "float4", "float8", "float16" }; |
| const char *sizeNames_double[] = { "double", "double2", "double4", "double8", "double16" }; |
| |
| test_status InitCL( cl_device_id device ) |
| { |
| int error; |
| uint32_t i, j; |
| int *bufSkip = NULL; |
| int isRTZ = 0; |
| RoundingMode oldRoundMode = kDefaultRoundingMode; |
| |
| 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; |
| |
| // check for cl_khr_fp64 |
| gHasDouble = is_extension_available(device, "cl_khr_fp64" ); |
| |
| if(0 == (CL_FP_INF_NAN & floatCapabilities) ) |
| gSkipNanInf = 1; |
| |
| // Embedded devices that flush to zero are allowed to have an undefined sign. |
| if (gIsEmbedded && gForceFTZ) |
| gIgnoreZeroSign = 1; |
| |
| gContext = clCreateContext( NULL, 1, &device, notify_callback, NULL, &error ); |
| if( NULL == gContext || error ) |
| { |
| vlog_error( "clCreateDeviceGroup failed. %d\n", error ); |
| return TEST_FAIL; |
| } |
| |
| gQueue = clCreateCommandQueue( gContext, device, 0, &error ); |
| if( NULL == gQueue || error ) |
| { |
| vlog_error( "clCreateContext failed. %d\n", error ); |
| return TEST_FAIL; |
| } |
| |
| // setup input buffers |
| bufA = clCreateBuffer( gContext, CL_MEM_READ_WRITE, BUFFER_SIZE, NULL, NULL ); |
| bufB = clCreateBuffer( gContext, CL_MEM_READ_WRITE, BUFFER_SIZE, NULL, NULL ); |
| bufC = clCreateBuffer( gContext, CL_MEM_READ_WRITE, BUFFER_SIZE, NULL, NULL ); |
| bufD = clCreateBuffer( gContext, CL_MEM_READ_WRITE, BUFFER_SIZE, NULL, NULL ); |
| bufE = clCreateBuffer( gContext, CL_MEM_READ_WRITE, BUFFER_SIZE, NULL, NULL ); |
| |
| if( bufA == NULL || |
| bufB == NULL || |
| bufC == NULL || |
| bufD == NULL || |
| bufE == NULL ) |
| { |
| vlog_error( "clCreateArray failed for input\n" ); |
| return TEST_FAIL; |
| } |
| |
| if( gHasDouble ) |
| { |
| bufC_double = clCreateBuffer( gContext, CL_MEM_READ_WRITE, BUFFER_SIZE, NULL, NULL ); |
| bufD_double = clCreateBuffer( gContext, CL_MEM_READ_WRITE, BUFFER_SIZE, NULL, NULL ); |
| if( bufC_double == NULL || |
| bufD_double == NULL ) |
| { |
| vlog_error( "clCreateArray failed for input DP\n" ); |
| return TEST_FAIL; |
| } |
| } |
| |
| const char *kernels[] = { |
| "", "#pragma OPENCL FP_CONTRACT OFF\n" |
| "__kernel void kernel1( __global ", NULL, " *out, const __global ", NULL, " *a, const __global ", NULL, " *b, const __global ", NULL, " *c )\n" |
| "{\n" |
| " int i = get_global_id(0);\n" |
| " out[i] = a[i] * b[i] + c[i];\n" |
| "}\n" |
| "\n" |
| "__kernel void kernel2( __global ", NULL, " *out, const __global ", NULL, " *a, const __global ", NULL, " *b, const __global ", NULL, " *c )\n" |
| "{\n" |
| " int i = get_global_id(0);\n" |
| " out[i] = a[i] * b[i] - c[i];\n" |
| "}\n" |
| "\n" |
| "__kernel void kernel3( __global ", NULL, " *out, const __global ", NULL, " *a, const __global ", NULL, " *b, const __global ", NULL, " *c )\n" |
| "{\n" |
| " int i = get_global_id(0);\n" |
| " out[i] = c[i] + a[i] * b[i];\n" |
| "}\n" |
| "\n" |
| "__kernel void kernel4( __global ", NULL, " *out, const __global ", NULL, " *a, const __global ", NULL, " *b, const __global ", NULL, " *c )\n" |
| "{\n" |
| " int i = get_global_id(0);\n" |
| " out[i] = c[i] - a[i] * b[i];\n" |
| "}\n" |
| "\n" |
| "__kernel void kernel5( __global ", NULL, " *out, const __global ", NULL, " *a, const __global ", NULL, " *b, const __global ", NULL, " *c )\n" |
| "{\n" |
| " int i = get_global_id(0);\n" |
| " out[i] = -(a[i] * b[i] + c[i]);\n" |
| "}\n" |
| "\n" |
| "__kernel void kernel6( __global ", NULL, " *out, const __global ", NULL, " *a, const __global ", NULL, " *b, const __global ", NULL, " *c )\n" |
| "{\n" |
| " int i = get_global_id(0);\n" |
| " out[i] = -(a[i] * b[i] - c[i]);\n" |
| "}\n" |
| "\n" |
| "__kernel void kernel7( __global ", NULL, " *out, const __global ", NULL, " *a, const __global ", NULL, " *b, const __global ", NULL, " *c )\n" |
| "{\n" |
| " int i = get_global_id(0);\n" |
| " out[i] = -(c[i] + a[i] * b[i]);\n" |
| "}\n" |
| "\n" |
| "__kernel void kernel8( __global ", NULL, " *out, const __global ", NULL, " *a, const __global ", NULL, " *b, const __global ", NULL, " *c )\n" |
| "{\n" |
| " int i = get_global_id(0);\n" |
| " out[i] = -(c[i] - a[i] * b[i]);\n" |
| "}\n" |
| "\n" }; |
| |
| for (i = 0; i < sizeof(sizeNames) / sizeof(sizeNames[0]); i++) |
| { |
| size_t strCount = sizeof(kernels) / sizeof(kernels[0]); |
| kernels[0] = ""; |
| |
| for (j = 2; j < strCount; j += 2) kernels[j] = sizeNames[i]; |
| error = create_single_kernel_helper(gContext, &gProgram[i], nullptr, |
| strCount, kernels, nullptr); |
| if (CL_SUCCESS != error || nullptr == gProgram[i]) |
| { |
| log_error("Error: Unable to create test program! (%s) (in %s:%d)\n", |
| IGetErrorString(error), __FILE__, __LINE__); |
| return TEST_FAIL; |
| } |
| } |
| |
| if (gHasDouble) |
| { |
| kernels[0] = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n"; |
| for (i = 0; i < sizeof(sizeNames_double) / sizeof(sizeNames_double[0]); |
| i++) |
| { |
| size_t strCount = sizeof(kernels) / sizeof(kernels[0]); |
| |
| for (j = 2; j < strCount; j += 2) kernels[j] = sizeNames_double[i]; |
| error = create_single_kernel_helper(gContext, &gProgram_double[i], |
| nullptr, strCount, kernels, |
| nullptr); |
| if (CL_SUCCESS != error || nullptr == gProgram_double[i]) |
| { |
| log_error( |
| "Error: Unable to create test program! (%s) (in %s:%d)\n", |
| IGetErrorString(error), __FILE__, __LINE__); |
| return TEST_FAIL; |
| } |
| } |
| } |
| |
| if( 0 == gSeedSpecified ) |
| { |
| time_t currentTime = time( NULL ); |
| struct tm *t = localtime(¤tTime); |
| gSeed = t->tm_sec + 60 * ( t->tm_min + 60 * (t->tm_hour + 24 * (t->tm_yday + 365 * t->tm_year))); |
| gSeed = (uint32_t) (((uint64_t) gSeed * (uint64_t) gSeed ) >> 16); |
| } |
| gMTdata = init_genrand( gSeed ); |
| |
| |
| // Init bufA and bufB |
| { |
| buf1 = (float *)malloc( BUFFER_SIZE ); |
| buf2 = (float *)malloc( BUFFER_SIZE ); |
| buf3 = (float *)malloc( BUFFER_SIZE ); |
| buf4 = (float *)malloc( BUFFER_SIZE ); |
| buf5 = (float *)malloc( BUFFER_SIZE ); |
| buf6 = (float *)malloc( BUFFER_SIZE ); |
| |
| bufSkip = (int *)malloc( BUFFER_SIZE ); |
| |
| if( NULL == buf1 || NULL == buf2 || NULL == buf3 || NULL == buf4 || NULL == buf5 || NULL == buf6 || NULL == bufSkip) |
| { |
| vlog_error( "Out of memory initializing buffers\n" ); |
| return TEST_FAIL; |
| } |
| for( i = 0; i < sizeof( correct ) / sizeof( correct[0] ); i++ ) |
| { |
| correct[i] = (float *)malloc( BUFFER_SIZE ); |
| skipTest[i] = (int *)malloc( BUFFER_SIZE ); |
| if(( NULL == correct[i] ) || ( NULL == skipTest[i])) |
| { |
| vlog_error( "Out of memory initializing buffers 2\n" ); |
| return TEST_FAIL; |
| } |
| } |
| |
| for( i = 0; i < BUFFER_SIZE / sizeof(float); i++ ) |
| ((uint32_t*) buf1)[i] = genrand_int32( gMTdata ); |
| |
| if( (error = clEnqueueWriteBuffer(gQueue, bufA, CL_FALSE, 0, BUFFER_SIZE, buf1, 0, NULL, NULL) )) |
| { |
| vlog_error( "Failure %d at clEnqueueWriteBuffer1\n", error ); |
| return TEST_FAIL; |
| } |
| |
| for( i = 0; i < BUFFER_SIZE / sizeof(float); i++ ) |
| ((uint32_t*) buf2)[i] = genrand_int32( gMTdata ); |
| |
| if( (error = clEnqueueWriteBuffer(gQueue, bufB, CL_FALSE, 0, BUFFER_SIZE, buf2, 0, NULL, NULL) )) |
| { |
| vlog_error( "Failure %d at clEnqueueWriteBuffer2\n", error ); |
| return TEST_FAIL; |
| } |
| |
| void *ftzInfo = NULL; |
| if( gForceFTZ ) |
| ftzInfo = FlushToZero(); |
| if ((CL_FP_ROUND_TO_ZERO == get_default_rounding_mode(device)) && gIsEmbedded) { |
| oldRoundMode = set_round(kRoundTowardZero, kfloat); |
| isRTZ = 1; |
| } |
| float *f = (float*) buf1; |
| float *f2 = (float*) buf2; |
| float *f3 = (float*) buf3; |
| float *f4 = (float*) buf4; |
| for( i = 0; i < BUFFER_SIZE / sizeof(float); i++ ) |
| { |
| float q = f[i]; |
| float q2 = f2[i]; |
| |
| feclearexcept(FE_OVERFLOW); |
| #if defined(_MSC_VER) && (defined(_M_IX86) || defined(_M_X64)) |
| // VS2005 might use x87 for straight multiplies, and we can't |
| // turn that off |
| f3[i] = sse_mul(q, q2); |
| f4[i] = sse_mul(-q, q2); |
| #elif defined(__PPC__) |
| // None of the current generation PPC processors support HW |
| // FTZ, emulate it in sw. |
| f3[i] = ppc_mul(q, q2); |
| f4[i] = ppc_mul(-q, q2); |
| #else |
| f3[i] = q * q2; |
| f4[i] = -q * q2; |
| #endif |
| // Skip test if the device doesn't support infinities and NaN AND the result overflows |
| // or either input is an infinity of NaN |
| bufSkip[i] = (gSkipNanInf && ((FE_OVERFLOW == (FE_OVERFLOW & fetestexcept(FE_OVERFLOW))) || |
| (fabsf(q) == FLT_MAX) || (q != q) || |
| (fabsf(q2) == FLT_MAX) || (q2 != q2))); |
| } |
| |
| if( gForceFTZ ) |
| UnFlushToZero(ftzInfo); |
| |
| if (isRTZ) |
| (void)set_round(oldRoundMode, kfloat); |
| |
| |
| if( (error = clEnqueueWriteBuffer(gQueue, bufC, CL_FALSE, 0, BUFFER_SIZE, buf3, 0, NULL, NULL) )) |
| { |
| vlog_error( "Failure %d at clEnqueueWriteBuffer3\n", error ); |
| return TEST_FAIL; |
| } |
| if( (error = clEnqueueWriteBuffer(gQueue, bufD, CL_FALSE, 0, BUFFER_SIZE, buf4, 0, NULL, NULL) )) |
| { |
| vlog_error( "Failure %d at clEnqueueWriteBuffer4\n", error ); |
| return TEST_FAIL; |
| } |
| |
| // Fill the buffers with NaN |
| float *f5 = (float*) buf5; |
| float nan_val = nanf(""); |
| for( i = 0; i < BUFFER_SIZE / sizeof( float ); i++ ) |
| f5[i] = nan_val; |
| |
| // calculate reference results |
| for( i = 0; i < BUFFER_SIZE / sizeof( float ); i++ ) |
| { |
| for ( j=0; j<8; j++) |
| { |
| feclearexcept(FE_OVERFLOW); |
| switch (j) |
| { |
| #if defined(_MSC_VER) && (defined(_M_IX86) || defined(_M_X64)) |
| // VS2005 might use x87 for straight add/sub, and we can't |
| // turn that off |
| case 0: |
| correct[0][i] = sse_add(buf3[i],buf4[i]); break; |
| case 1: |
| correct[1][i] = sse_sub(buf3[i],buf3[i]); break; |
| case 2: |
| correct[2][i] = sse_add(buf4[i],buf3[i]); break; |
| case 3: |
| correct[3][i] = sse_sub(buf3[i],buf3[i]); break; |
| case 4: |
| correct[4][i] = -sse_add(buf3[i],buf4[i]); break; |
| case 5: |
| correct[5][i] = -sse_sub(buf3[i],buf3[i]); break; |
| case 6: |
| correct[6][i] = -sse_add(buf4[i],buf3[i]); break; |
| case 7: |
| correct[7][i] = -sse_sub(buf3[i],buf3[i]); break; |
| #else |
| case 0: |
| correct[0][i] = buf3[i] + buf4[i]; break; |
| case 1: |
| correct[1][i] = buf3[i] - buf3[i]; break; |
| case 2: |
| correct[2][i] = buf4[i] + buf3[i]; break; |
| case 3: |
| correct[3][i] = buf3[i] - buf3[i]; break; |
| case 4: |
| correct[4][i] = -(buf3[i] + buf4[i]); break; |
| case 5: |
| correct[5][i] = -(buf3[i] - buf3[i]); break; |
| case 6: |
| correct[6][i] = -(buf4[i] + buf3[i]); break; |
| case 7: |
| correct[7][i] = -(buf3[i] - buf3[i]); break; |
| #endif |
| } |
| // Further skip test inputs if the device doesn support infinities AND NaNs |
| // resulting sum overflows |
| skipTest[j][i] = (bufSkip[i] || |
| (gSkipNanInf && (FE_OVERFLOW == (FE_OVERFLOW & fetestexcept(FE_OVERFLOW))))); |
| |
| #if defined(__PPC__) |
| // Since the current Power processors don't emulate flush to zero in HW, |
| // it must be emulated in SW instead. |
| if (gForceFTZ) |
| { |
| if ((fabsf(correct[j][i]) < FLT_MIN) && (correct[j][i] != 0.0f)) |
| correct[j][i] = copysignf(0.0f, correct[j][i]); |
| } |
| #endif |
| } |
| } |
| if( gHasDouble ) |
| { |
| // Spec requires correct non-flushed results |
| // for doubles. We disable FTZ if this is default on |
| // the platform (like ARM) for reference result computation |
| // It is no-op if platform default is not FTZ (e.g. x86) |
| FPU_mode_type oldMode; |
| DisableFTZ( &oldMode ); |
| |
| buf3_double = (double *)malloc( BUFFER_SIZE ); |
| buf4_double = (double *)malloc( BUFFER_SIZE ); |
| buf5_double = (double *)malloc( BUFFER_SIZE ); |
| buf6_double = (double *)malloc( BUFFER_SIZE ); |
| if( NULL == buf3_double || NULL == buf4_double || NULL == buf5_double || NULL == buf6_double ) |
| { |
| vlog_error( "Out of memory initializing DP buffers\n" ); |
| return TEST_FAIL; |
| } |
| for( i = 0; i < sizeof( correct_double ) / sizeof( correct_double[0] ); i++ ) |
| { |
| correct_double[i] = (double *)malloc( BUFFER_SIZE ); |
| if( NULL == correct_double[i] ) |
| { |
| vlog_error( "Out of memory initializing DP buffers 2\n" ); |
| return TEST_FAIL; |
| } |
| } |
| |
| |
| double *f = (double*) buf1; |
| double *f2 = (double*) buf2; |
| double *f3 = (double*) buf3_double; |
| double *f4 = (double*) buf4_double; |
| for( i = 0; i < BUFFER_SIZE / sizeof(double); i++ ) |
| { |
| double q = f[i]; |
| double q2 = f2[i]; |
| #if defined(_MSC_VER) && (defined(_M_IX86) || defined(_M_X64)) |
| // VS2005 might use x87 for straight multiplies, and we can't |
| // turn that off |
| f3[i] = sse_mul_sd(q, q2); |
| f4[i] = sse_mul_sd(-q, q2); |
| #else |
| f3[i] = q * q2; |
| f4[i] = -q * q2; |
| #endif |
| } |
| |
| if( (error = clEnqueueWriteBuffer(gQueue, bufC_double, CL_FALSE, 0, BUFFER_SIZE, buf3_double, 0, NULL, NULL) )) |
| { |
| vlog_error( "Failure %d at clEnqueueWriteBuffer3\n", error ); |
| return TEST_FAIL; |
| } |
| if( (error = clEnqueueWriteBuffer(gQueue, bufD_double, CL_FALSE, 0, BUFFER_SIZE, buf4_double, 0, NULL, NULL) )) |
| { |
| vlog_error( "Failure %d at clEnqueueWriteBuffer4\n", error ); |
| return TEST_FAIL; |
| } |
| |
| // Fill the buffers with NaN |
| double *f5 = (double*) buf5_double; |
| double nan_val = nanf(""); |
| for( i = 0; i < BUFFER_SIZE / sizeof( double ); i++ ) |
| f5[i] = nan_val; |
| |
| // calculate reference results |
| for( i = 0; i < BUFFER_SIZE / sizeof( double ); i++ ) |
| { |
| #if defined(_MSC_VER) && (defined(_M_IX86) || defined(_M_X64)) |
| // VS2005 might use x87 for straight add/sub, and we can't |
| // turn that off |
| correct_double[0][i] = sse_add_sd(buf3_double[i],buf4_double[i]); |
| correct_double[1][i] = sse_sub_sd(buf3_double[i],buf3_double[i]); |
| correct_double[2][i] = sse_add_sd(buf4_double[i],buf3_double[i]); |
| correct_double[3][i] = sse_sub_sd(buf3_double[i],buf3_double[i]); |
| correct_double[4][i] = -sse_add_sd(buf3_double[i],buf4_double[i]); |
| correct_double[5][i] = -sse_sub_sd(buf3_double[i],buf3_double[i]); |
| correct_double[6][i] = -sse_add_sd(buf4_double[i],buf3_double[i]); |
| correct_double[7][i] = -sse_sub_sd(buf3_double[i],buf3_double[i]); |
| #else |
| correct_double[0][i] = buf3_double[i] + buf4_double[i]; |
| correct_double[1][i] = buf3_double[i] - buf3_double[i]; |
| correct_double[2][i] = buf4_double[i] + buf3_double[i]; |
| correct_double[3][i] = buf3_double[i] - buf3_double[i]; |
| correct_double[4][i] = -(buf3_double[i] + buf4_double[i]); |
| correct_double[5][i] = -(buf3_double[i] - buf3_double[i]); |
| correct_double[6][i] = -(buf4_double[i] + buf3_double[i]); |
| correct_double[7][i] = -(buf3_double[i] - buf3_double[i]); |
| #endif |
| } |
| |
| // Restore previous FP state since we modified it for |
| // reference result computation (see DisableFTZ call above) |
| RestoreFPState(&oldMode); |
| } |
| } |
| |
| char c[1000]; |
| static const char *no_yes[] = { "NO", "YES" }; |
| vlog( "\nCompute Device info:\n" ); |
| clGetDeviceInfo( device, CL_DEVICE_NAME, sizeof(c), (void *)&c, NULL); |
| vlog( "\tDevice Name: %s\n", c ); |
| clGetDeviceInfo( device, CL_DEVICE_VENDOR, sizeof(c), (void *)&c, NULL); |
| vlog( "\tVendor: %s\n", c ); |
| clGetDeviceInfo( device, CL_DEVICE_VERSION, sizeof(c), (void *)&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), (void *)&c, NULL); |
| vlog( "\tDriver Version: %s\n", c ); |
| vlog( "\tSubnormal values supported? %s\n", no_yes[0 != (CL_FP_DENORM & floatCapabilities)] ); |
| vlog( "\tTesting with FTZ mode ON? %s\n", no_yes[0 != gForceFTZ] ); |
| vlog( "\tTesting Doubles? %s\n", no_yes[0 != gHasDouble] ); |
| vlog( "\tRandom Number seed: 0x%8.8x\n", gSeed ); |
| vlog( "\n\n" ); |
| |
| return TEST_PASS; |
| } |
| |
| static void ReleaseCL( void ) |
| { |
| clReleaseMemObject(bufA); |
| clReleaseMemObject(bufB); |
| clReleaseMemObject(bufC); |
| clReleaseMemObject(bufD); |
| clReleaseMemObject(bufE); |
| clReleaseProgram(gProgram[0]); |
| clReleaseProgram(gProgram[1]); |
| clReleaseProgram(gProgram[2]); |
| clReleaseProgram(gProgram[3]); |
| clReleaseProgram(gProgram[4]); |
| if( gHasDouble ) |
| { |
| clReleaseMemObject(bufC_double); |
| clReleaseMemObject(bufD_double); |
| clReleaseProgram(gProgram_double[0]); |
| clReleaseProgram(gProgram_double[1]); |
| clReleaseProgram(gProgram_double[2]); |
| clReleaseProgram(gProgram_double[3]); |
| clReleaseProgram(gProgram_double[4]); |
| } |
| clReleaseCommandQueue(gQueue); |
| clReleaseContext(gContext); |
| } |
| |
| |
| static int RunTest( int testNumber ) |
| { |
| size_t i; |
| int error = 0; |
| cl_mem args[4]; |
| float *c; |
| const char *kernelName[] = { "kernel1", "kernel2", "kernel3", "kernel4", |
| "kernel5", "kernel6", "kernel7", "kernel8" }; |
| switch( testNumber ) |
| { |
| case 0: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufD; c = buf4; break; // a * b + c |
| case 1: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufC; c = buf3; break; |
| case 2: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufD; c = buf4; break; |
| case 3: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufC; c = buf3; break; |
| case 4: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufD; c = buf4; break; |
| case 5: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufC; c = buf3; break; |
| case 6: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufD; c = buf4; break; |
| case 7: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufC; c = buf3; break; |
| default: |
| vlog_error( "Unknown test case %d passed to RunTest\n", testNumber ); |
| return -1; |
| } |
| |
| |
| int vectorSize; |
| for( vectorSize = 0; vectorSize < 5; vectorSize++ ) |
| { |
| cl_kernel k = clCreateKernel( gProgram[ vectorSize ], kernelName[ testNumber ], &error ); |
| if( NULL == k || error ) |
| { |
| vlog_error( "%d) Unable to find kernel \"%s\" for vector size: %d\n", error, kernelName[ testNumber ], 1 << vectorSize ); |
| return -2; |
| } |
| |
| // set the kernel args |
| for( i = 0; i < sizeof(args ) / sizeof( args[0]); i++ ) |
| if( (error = clSetKernelArg(k, i, sizeof( cl_mem ), args + i) )) |
| { |
| vlog_error( "Error %d setting kernel arg # %ld\n", error, i ); |
| return error; |
| } |
| |
| // write NaNs to the result array |
| if( (error = clEnqueueWriteBuffer(gQueue, bufE, CL_TRUE, 0, BUFFER_SIZE, buf5, 0, NULL, NULL) )) |
| { |
| vlog_error( "Failure %d at clWriteArray %d\n", error, testNumber ); |
| return error; |
| } |
| |
| // execute the kernel |
| size_t gDim[3] = { BUFFER_SIZE / (sizeof( cl_float ) * (1<<vectorSize)), 0, 0 }; |
| if( ((error = clEnqueueNDRangeKernel(gQueue, k, 1, NULL, gDim, NULL, 0, NULL, NULL) ))) |
| { |
| vlog_error( "Got Error # %d trying to execture kernel\n", error ); |
| return error; |
| } |
| |
| // read the data back |
| if( (error = clEnqueueReadBuffer(gQueue, bufE, CL_TRUE, 0, BUFFER_SIZE, buf6, 0, NULL, NULL ) )) |
| { |
| vlog_error( "Failure %d at clReadArray %d\n", error, testNumber ); |
| return error; |
| } |
| |
| // verify results |
| float *test = (float*) buf6; |
| float *a = (float*) buf1; |
| float *b = (float*) buf2; |
| for( i = 0; i < BUFFER_SIZE / sizeof( float ); i++ ) |
| { |
| if( isnan(test[i]) && isnan(correct[testNumber][i] ) ) |
| continue; |
| |
| if( skipTest[testNumber][i] ) |
| continue; |
| |
| // sign of zero must be correct |
| if(( ((uint32_t*) test)[i] != ((uint32_t*) correct[testNumber])[i] ) && |
| !(gIgnoreZeroSign && (test[i] == 0.0f) && (correct[testNumber][i] == 0.0f)) ) |
| { |
| switch( testNumber ) |
| { |
| // Zeros for these should be positive |
| case 0: vlog_error( "%ld) Error for %s %s: %a * %a + %a = *%a vs. %a\n", i, sizeNames[ vectorSize], kernelName[ testNumber ], |
| a[i], b[i], c[i], correct[testNumber][i], test[i] ); clReleaseKernel(k); return -1; |
| case 1: vlog_error( "%ld) Error for %s %s: %a * %a - %a = *%a vs. %a\n", i, sizeNames[ vectorSize], kernelName[ testNumber ], |
| a[i], b[i], c[i], correct[testNumber][i], test[i] ); clReleaseKernel(k); return -1; |
| case 2: vlog_error( "%ld) Error for %s %s: %a + %a * %a = *%a vs. %a\n", i, sizeNames[ vectorSize], kernelName[ testNumber ], |
| c[i], a[i], b[i], correct[testNumber][i], test[i] ); clReleaseKernel(k); return -1; |
| case 3: vlog_error( "%ld) Error for %s %s: %a - %a * %a = *%a vs. %a\n", i, sizeNames[ vectorSize], kernelName[ testNumber ], |
| c[i], a[i], b[i], correct[testNumber][i], test[i] ); clReleaseKernel(k); return -1; |
| |
| // Zeros for these should be negative |
| case 4: vlog_error( "%ld) Error for %s %s: -(%a * %a + %a) = *%a vs. %a\n", i, sizeNames[ vectorSize], kernelName[ testNumber ], |
| a[i], b[i], c[i], correct[testNumber][i], test[i] ); clReleaseKernel(k); return -1; |
| case 5: vlog_error( "%ld) Error for %s %s: -(%a * %a - %a) = *%a vs. %a\n", i, sizeNames[ vectorSize], kernelName[ testNumber ], |
| a[i], b[i], c[i], correct[testNumber][i], test[i] ); clReleaseKernel(k); return -1; |
| case 6: vlog_error( "%ld) Error for %s %s: -(%a + %a * %a) = *%a vs. %a\n", i, sizeNames[ vectorSize], kernelName[ testNumber ], |
| c[i], a[i], b[i], correct[testNumber][i], test[i] ); clReleaseKernel(k); return -1; |
| case 7: vlog_error( "%ld) Error for %s %s: -(%a - %a * %a) = *%a vs. %a\n", i, sizeNames[ vectorSize], kernelName[ testNumber ], |
| c[i], a[i], b[i], correct[testNumber][i], test[i] ); clReleaseKernel(k); return -1; |
| default: |
| vlog_error( "error: Unknown test number!\n" ); |
| clReleaseKernel(k); |
| return -2; |
| } |
| } |
| } |
| |
| clReleaseKernel(k); |
| } |
| |
| return error; |
| } |
| |
| static int RunTest_Double( int testNumber ) |
| { |
| if( !gHasDouble ) |
| { |
| vlog("Double is not supported, test not run.\n"); |
| return 0; |
| } |
| |
| size_t i; |
| int error = 0; |
| cl_mem args[4]; |
| double *c; |
| const char *kernelName[] = { "kernel1", "kernel2", "kernel3", "kernel4", |
| "kernel5", "kernel6", "kernel7", "kernel8" }; |
| |
| switch( testNumber ) |
| { |
| case 0: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufD_double; c = buf4_double; break; // a * b + c |
| case 1: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufC_double; c = buf3_double; break; |
| case 2: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufD_double; c = buf4_double; break; |
| case 3: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufC_double; c = buf3_double; break; |
| case 4: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufD_double; c = buf4_double; break; |
| case 5: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufC_double; c = buf3_double; break; |
| case 6: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufD_double; c = buf4_double; break; |
| case 7: args[0] = bufE; args[1] = bufA; args[2] = bufB; args[3] = bufC_double; c = buf3_double; break; |
| default: |
| vlog_error( "Unknown test case %d passed to RunTest\n", testNumber ); |
| return -1; |
| } |
| |
| int vectorSize; |
| for( vectorSize = 0; vectorSize < 5; vectorSize++ ) |
| { |
| cl_kernel k = clCreateKernel( gProgram_double[ vectorSize ], kernelName[ testNumber ], &error ); |
| if( NULL == k || error ) |
| { |
| vlog_error( "%d) Unable to find kernel \"%s\" for vector size: %d\n", error, kernelName[ testNumber ], 1 << vectorSize ); |
| return -2; |
| } |
| |
| // set the kernel args |
| for( i = 0; i < sizeof(args ) / sizeof( args[0]); i++ ) |
| if( (error = clSetKernelArg(k, i, sizeof( cl_mem ), args + i) )) |
| { |
| vlog_error( "Error %d setting kernel arg # %ld\n", error, i ); |
| return error; |
| } |
| |
| // write NaNs to the result array |
| if( (error = clEnqueueWriteBuffer(gQueue, bufE, CL_FALSE, 0, BUFFER_SIZE, buf5_double, 0, NULL, NULL) )) |
| { |
| vlog_error( "Failure %d at clWriteArray %d\n", error, testNumber ); |
| return error; |
| } |
| |
| // execute the kernel |
| size_t gDim[3] = { BUFFER_SIZE / (sizeof( cl_double ) * (1<<vectorSize)), 0, 0 }; |
| if( ((error = clEnqueueNDRangeKernel(gQueue, k, 1, NULL, gDim, NULL, 0, NULL, NULL) ))) |
| { |
| vlog_error( "Got Error # %d trying to execture kernel\n", error ); |
| return error; |
| } |
| |
| // read the data back |
| if( (error = clEnqueueReadBuffer(gQueue, bufE, CL_TRUE, 0, BUFFER_SIZE, buf6_double, 0, NULL, NULL ) )) |
| { |
| vlog_error( "Failure %d at clReadArray %d\n", error, testNumber ); |
| return error; |
| } |
| |
| // verify results |
| double *test = (double*) buf6_double; |
| double *a = (double*) buf1; |
| double *b = (double*) buf2; |
| for( i = 0; i < BUFFER_SIZE / sizeof( double ); i++ ) |
| { |
| if( isnan(test[i]) && isnan(correct_double[testNumber][i] ) ) |
| continue; |
| |
| // sign of zero must be correct |
| if( ((uint64_t*) test)[i] != ((uint64_t*) correct_double[testNumber])[i] ) |
| { |
| switch( testNumber ) |
| { |
| // Zeros for these should be positive |
| case 0: vlog_error( "%ld) Error for %s %s: %a * %a + %a = *%a vs. %a\n", i, sizeNames_double[ vectorSize], kernelName[ testNumber ], |
| a[i], b[i], c[i], correct[testNumber][i], test[i] ); return -1; |
| case 1: vlog_error( "%ld) Error for %s %s: %a * %a - %a = *%a vs. %a\n", i, sizeNames_double[ vectorSize], kernelName[ testNumber ], |
| a[i], b[i], c[i], correct[testNumber][i], test[i] ); return -1; |
| case 2: vlog_error( "%ld) Error for %s %s: %a + %a * %a = *%a vs. %a\n", i, sizeNames_double[ vectorSize], kernelName[ testNumber ], |
| c[i], a[i], b[i], correct[testNumber][i], test[i] ); return -1; |
| case 3: vlog_error( "%ld) Error for %s %s: %a - %a * %a = *%a vs. %a\n", i, sizeNames_double[ vectorSize], kernelName[ testNumber ], |
| c[i], a[i], b[i], correct[testNumber][i], test[i] ); return -1; |
| |
| // Zeros for these should be negative |
| case 4: vlog_error( "%ld) Error for %s %s: -(%a * %a + %a) = *%a vs. %a\n", i, sizeNames_double[ vectorSize], kernelName[ testNumber ], |
| a[i], b[i], c[i], correct[testNumber][i], test[i] ); return -1; |
| case 5: vlog_error( "%ld) Error for %s %s: -(%a * %a - %a) = *%a vs. %a\n", i, sizeNames_double[ vectorSize], kernelName[ testNumber ], |
| a[i], b[i], c[i], correct[testNumber][i], test[i] ); return -1; |
| case 6: vlog_error( "%ld) Error for %s %s: -(%a + %a * %a) = *%a vs. %a\n", i, sizeNames_double[ vectorSize], kernelName[ testNumber ], |
| c[i], a[i], b[i], correct[testNumber][i], test[i] ); return -1; |
| case 7: vlog_error( "%ld) Error for %s %s: -(%a - %a * %a) = *%a vs. %a\n", i, sizeNames_double[ vectorSize], kernelName[ testNumber ], |
| c[i], a[i], b[i], correct[testNumber][i], test[i] ); return -1; |
| default: |
| vlog_error( "error: Unknown test number!\n" ); |
| return -2; |
| } |
| } |
| } |
| |
| clReleaseKernel(k); |
| } |
| |
| return error; |
| } |
