// // 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 "testBase.h" #include "harness/conversions.h" #include "harness/typeWrappers.h" #include "harness/testHarness.h" const char *anyAllTestKernelPattern = "%s\n" // optional pragma "__kernel void sample_test(__global %s%s *sourceA, __global int *destValues)\n" "{\n" " int tid = get_global_id(0);\n" " destValues[tid] = %s( sourceA[tid] );\n" "\n" "}\n"; const char *anyAllTestKernelPatternVload = "%s\n" // optional pragma "__kernel void sample_test(__global %s%s *sourceA, __global int *destValues)\n" "{\n" " int tid = get_global_id(0);\n" " destValues[tid] = %s(vload3(tid, (__global %s *)sourceA));\n" // ugh, almost "\n" "}\n"; #define TEST_SIZE 512 typedef int (*anyAllVerifyFn)( ExplicitType vecType, unsigned int vecSize, void *inData ); int test_any_all_kernel(cl_context context, cl_command_queue queue, const char *fnName, ExplicitType vecType, unsigned int vecSize, anyAllVerifyFn verifyFn, MTdata d ) { clProgramWrapper program; clKernelWrapper kernel; clMemWrapper streams[2]; cl_long inDataA[TEST_SIZE * 16], clearData[TEST_SIZE * 16]; int outData[TEST_SIZE]; int error, i; size_t threads[1], localThreads[1]; char kernelSource[10240]; char *programPtr; char sizeName[4]; /* Create the source */ if( g_vector_aligns[vecSize] == 1 ) { sizeName[ 0 ] = 0; } else { sprintf( sizeName, "%d", vecSize ); } log_info("Testing any/all on %s%s\n", get_explicit_type_name( vecType ), sizeName); if(DENSE_PACK_VECS && vecSize == 3) { // anyAllTestKernelPatternVload sprintf( kernelSource, anyAllTestKernelPatternVload, vecType == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : "", get_explicit_type_name( vecType ), sizeName, fnName, get_explicit_type_name(vecType)); } else { sprintf( kernelSource, anyAllTestKernelPattern, vecType == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : "", get_explicit_type_name( vecType ), sizeName, fnName ); } /* Create kernels */ programPtr = kernelSource; if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "sample_test" ) ) { return -1; } /* Generate some streams */ generate_random_data( vecType, TEST_SIZE * g_vector_aligns[vecSize], d, inDataA ); memset( clearData, 0, sizeof( clearData ) ); streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, get_explicit_type_size(vecType) * g_vector_aligns[vecSize] * TEST_SIZE, &inDataA, &error); if( streams[0] == NULL ) { print_error( error, "Creating input array A failed!\n"); return -1; } streams[1] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(cl_int) * g_vector_aligns[vecSize] * TEST_SIZE, clearData, &error); if( streams[1] == NULL ) { print_error( error, "Creating output array failed!\n"); return -1; } /* Assign streams and execute */ error = clSetKernelArg( kernel, 0, sizeof( streams[0] ), &streams[0] ); test_error( error, "Unable to set indexed kernel arguments" ); error = clSetKernelArg( kernel, 1, sizeof( streams[1] ), &streams[1] ); test_error( error, "Unable to set indexed kernel arguments" ); /* Run the kernel */ threads[0] = TEST_SIZE; error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] ); test_error( error, "Unable to get work group size to use" ); error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL ); test_error( error, "Unable to execute test kernel" ); /* Now get the results */ error = clEnqueueReadBuffer( queue, streams[1], true, 0, sizeof( int ) * TEST_SIZE, outData, 0, NULL, NULL ); test_error( error, "Unable to read output array!" ); /* And verify! */ for( i = 0; i < TEST_SIZE; i++ ) { int expected = verifyFn( vecType, vecSize, (char *)inDataA + i * get_explicit_type_size( vecType ) * g_vector_aligns[vecSize] ); if( expected != outData[ i ] ) { unsigned int *ptr = (unsigned int *)( (char *)inDataA + i * get_explicit_type_size( vecType ) * g_vector_aligns[vecSize] ); log_error( "ERROR: Data sample %d does not validate! Expected (%d), got (%d), source 0x%08x\n", i, expected, outData[i], *ptr ); return -1; } } return 0; } int anyVerifyFn( ExplicitType vecType, unsigned int vecSize, void *inData ) { unsigned int i; switch( vecType ) { case kChar: { char sum = 0; char *tData = (char *)inData; for( i = 0; i < vecSize; i++ ) sum |= tData[ i ] & 0x80; return (sum != 0) ? 1 : 0; } case kShort: { short sum = 0; short *tData = (short *)inData; for( i = 0; i < vecSize; i++ ) sum |= tData[ i ] & 0x8000; return (sum != 0); } case kInt: { cl_int sum = 0; cl_int *tData = (cl_int *)inData; for( i = 0; i < vecSize; i++ ) sum |= tData[ i ] & (cl_int)0x80000000L; return (sum != 0); } case kLong: { cl_long sum = 0; cl_long *tData = (cl_long *)inData; for( i = 0; i < vecSize; i++ ) sum |= tData[ i ] & 0x8000000000000000LL; return (sum != 0); } default: return 0; } } int test_relational_any(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { ExplicitType vecType[] = { kChar, kShort, kInt, kLong }; unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 }; unsigned int index, typeIndex; int retVal = 0; RandomSeed seed(gRandomSeed ); for( typeIndex = 0; typeIndex < 4; typeIndex++ ) { if (vecType[typeIndex] == kLong && !gHasLong) continue; for( index = 0; vecSizes[ index ] != 0; index++ ) { // Test! if( test_any_all_kernel(context, queue, "any", vecType[ typeIndex ], vecSizes[ index ], anyVerifyFn, seed ) != 0 ) { log_error( " Vector %s%d FAILED\n", get_explicit_type_name( vecType[ typeIndex ] ), vecSizes[ index ] ); retVal = -1; } } } return retVal; } int allVerifyFn( ExplicitType vecType, unsigned int vecSize, void *inData ) { unsigned int i; switch( vecType ) { case kChar: { char sum = 0x80; char *tData = (char *)inData; for( i = 0; i < vecSize; i++ ) sum &= tData[ i ] & 0x80; return (sum != 0) ? 1 : 0; } case kShort: { short sum = 0x8000; short *tData = (short *)inData; for( i = 0; i < vecSize; i++ ) sum &= tData[ i ] & 0x8000; return (sum != 0); } case kInt: { cl_int sum = 0x80000000L; cl_int *tData = (cl_int *)inData; for( i = 0; i < vecSize; i++ ) sum &= tData[ i ] & (cl_int)0x80000000L; return (sum != 0); } case kLong: { cl_long sum = 0x8000000000000000LL; cl_long *tData = (cl_long *)inData; for( i = 0; i < vecSize; i++ ) sum &= tData[ i ] & 0x8000000000000000LL; return (sum != 0); } default: return 0; } } int test_relational_all(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { ExplicitType vecType[] = { kChar, kShort, kInt, kLong }; unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 }; unsigned int index, typeIndex; int retVal = 0; RandomSeed seed(gRandomSeed ); for( typeIndex = 0; typeIndex < 4; typeIndex++ ) { if (vecType[typeIndex] == kLong && !gHasLong) continue; for( index = 0; vecSizes[ index ] != 0; index++ ) { // Test! if( test_any_all_kernel(context, queue, "all", vecType[ typeIndex ], vecSizes[ index ], allVerifyFn, seed ) != 0 ) { log_error( " Vector %s%d FAILED\n", get_explicit_type_name( vecType[ typeIndex ] ), vecSizes[ index ] ); retVal = -1; } } } return retVal; } const char *selectTestKernelPattern = "%s\n" // optional pragma "__kernel void sample_test(__global %s%s *sourceA, __global %s%s *sourceB, __global %s%s *sourceC, __global %s%s *destValues)\n" "{\n" " int tid = get_global_id(0);\n" " destValues[tid] = %s( sourceA[tid], sourceB[tid], sourceC[tid] );\n" "\n" "}\n"; const char *selectTestKernelPatternVload = "%s\n" // optional pragma "__kernel void sample_test(__global %s%s *sourceA, __global %s%s *sourceB, __global %s%s *sourceC, __global %s%s *destValues)\n" "{\n" " int tid = get_global_id(0);\n" " %s%s tmp = %s( vload3(tid, (__global %s *)sourceA), vload3(tid, (__global %s *)sourceB), vload3(tid, (__global %s *)sourceC) );\n" " vstore3(tmp, tid, (__global %s *)destValues);\n" "\n" "}\n"; typedef void (*selectVerifyFn)( ExplicitType vecType, ExplicitType testVecType, unsigned int vecSize, void *inDataA, void *inDataB, void *inDataTest, void *outData ); int test_select_kernel(cl_context context, cl_command_queue queue, const char *fnName, ExplicitType vecType, unsigned int vecSize, ExplicitType testVecType, selectVerifyFn verifyFn, MTdata d ) { clProgramWrapper program; clKernelWrapper kernel; clMemWrapper streams[4]; cl_long inDataA[TEST_SIZE * 16], inDataB[ TEST_SIZE * 16 ], inDataC[ TEST_SIZE * 16 ]; cl_long outData[TEST_SIZE * 16], expected[16]; int error, i; size_t threads[1], localThreads[1]; char kernelSource[10240]; char *programPtr; char sizeName[4], outSizeName[4]; unsigned int outVecSize; /* Create the source */ if( vecSize == 1 ) sizeName[ 0 ] = 0; else sprintf( sizeName, "%d", vecSize ); outVecSize = vecSize; if( outVecSize == 1 ) outSizeName[ 0 ] = 0; else sprintf( outSizeName, "%d", outVecSize ); if(DENSE_PACK_VECS && vecSize == 3) { // anyAllTestKernelPatternVload sprintf( kernelSource, selectTestKernelPatternVload, (vecType == kDouble || testVecType == kDouble) ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : "", get_explicit_type_name( vecType ), sizeName, get_explicit_type_name( vecType ), sizeName, get_explicit_type_name( testVecType ), sizeName, get_explicit_type_name( vecType ), outSizeName, get_explicit_type_name( vecType ), sizeName, fnName, get_explicit_type_name( vecType ), get_explicit_type_name( vecType ), get_explicit_type_name( vecType ), get_explicit_type_name( testVecType ) ); } else { sprintf( kernelSource, selectTestKernelPattern, (vecType == kDouble || testVecType == kDouble) ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : "", get_explicit_type_name( vecType ), sizeName, get_explicit_type_name( vecType ), sizeName, get_explicit_type_name( testVecType ), sizeName, get_explicit_type_name( vecType ), outSizeName, fnName ); } /* Create kernels */ programPtr = kernelSource; if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "sample_test" ) ) { return -1; } /* Generate some streams */ generate_random_data( vecType, TEST_SIZE * g_vector_aligns[vecSize], d, inDataA ); generate_random_data( vecType, TEST_SIZE * g_vector_aligns[vecSize], d, inDataB ); generate_random_data( testVecType, TEST_SIZE * g_vector_aligns[vecSize], d, inDataC ); streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, get_explicit_type_size(vecType) * g_vector_aligns[vecSize] * TEST_SIZE, &inDataA, &error); if( streams[0] == NULL ) { print_error( error, "Creating input array A failed!\n"); return -1; } streams[1] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, get_explicit_type_size(vecType) * g_vector_aligns[vecSize] * TEST_SIZE, &inDataB, &error); if( streams[1] == NULL ) { print_error( error, "Creating input array A failed!\n"); return -1; } streams[2] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, get_explicit_type_size(testVecType) * g_vector_aligns[vecSize] * TEST_SIZE, &inDataC, &error); if( streams[2] == NULL ) { print_error( error, "Creating input array A failed!\n"); return -1; } streams[3] = clCreateBuffer( context, CL_MEM_READ_WRITE, get_explicit_type_size( vecType ) * g_vector_aligns[outVecSize] * TEST_SIZE, NULL, &error); if( streams[3] == NULL ) { print_error( error, "Creating output array failed!\n"); return -1; } /* Assign streams and execute */ error = clSetKernelArg( kernel, 0, sizeof( streams[0] ), &streams[0] ); test_error( error, "Unable to set indexed kernel arguments" ); error = clSetKernelArg( kernel, 1, sizeof( streams[1] ), &streams[1] ); test_error( error, "Unable to set indexed kernel arguments" ); error = clSetKernelArg( kernel, 2, sizeof( streams[2] ), &streams[2] ); test_error( error, "Unable to set indexed kernel arguments" ); error = clSetKernelArg( kernel, 3, sizeof( streams[3] ), &streams[3] ); test_error( error, "Unable to set indexed kernel arguments" ); /* Run the kernel */ threads[0] = TEST_SIZE; error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] ); test_error( error, "Unable to get work group size to use" ); error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL ); test_error( error, "Unable to execute test kernel" ); /* Now get the results */ error = clEnqueueReadBuffer( queue, streams[3], true, 0, get_explicit_type_size( vecType ) * TEST_SIZE * g_vector_aligns[outVecSize], outData, 0, NULL, NULL ); test_error( error, "Unable to read output array!" ); /* And verify! */ for( i = 0; i < (int)(TEST_SIZE * g_vector_aligns[vecSize]); i++ ) { if(i%g_vector_aligns[vecSize] >= (int) vecSize) { continue; } verifyFn( vecType, testVecType, vecSize, (char *)inDataA + i * get_explicit_type_size( vecType ), (char *)inDataB + i * get_explicit_type_size( vecType ), (char *)inDataC + i * get_explicit_type_size( testVecType ), expected); char *outPtr = (char *)outData; outPtr += ( i / g_vector_aligns[vecSize] ) * get_explicit_type_size( vecType ) * g_vector_aligns[outVecSize]; outPtr += ( i % g_vector_aligns[vecSize] ) * get_explicit_type_size( vecType ); if( memcmp( expected, outPtr, get_explicit_type_size( vecType ) ) != 0 ) { log_error( "ERROR: Data sample %d:%d does not validate! Expected (0x%08x), got (0x%08x) from (0x%08x) and (0x%08x) with test (0x%08x)\n", i / g_vector_aligns[vecSize], i % g_vector_aligns[vecSize], *( (int *)expected ), *( (int *)( (char *)outData + i * get_explicit_type_size( vecType ) ) ), *( (int *)( (char *)inDataA + i * get_explicit_type_size( vecType ) ) ), *( (int *)( (char *)inDataB + i * get_explicit_type_size( vecType ) ) ), *( (int *)( (char *)inDataC + i*get_explicit_type_size( testVecType ) ) ) ); int j; log_error( "inA: " ); unsigned char *a = (unsigned char *)( (char *)inDataA + i * get_explicit_type_size( vecType ) ); unsigned char *b = (unsigned char *)( (char *)inDataB + i * get_explicit_type_size( vecType ) ); unsigned char *c = (unsigned char *)( (char *)inDataC + i * get_explicit_type_size( testVecType ) ); unsigned char *e = (unsigned char *)( expected ); unsigned char *g = (unsigned char *)( (char *)outData + i * get_explicit_type_size( vecType ) ); for( j = 0; j < 16; j++ ) log_error( "0x%02x ", a[ j ] ); log_error( "\ninB: " ); for( j = 0; j < 16; j++ ) log_error( "0x%02x ", b[ j ] ); log_error( "\ninC: " ); for( j = 0; j < 16; j++ ) log_error( "0x%02x ", c[ j ] ); log_error( "\nexp: " ); for( j = 0; j < 16; j++ ) log_error( "0x%02x ", e[ j ] ); log_error( "\ngot: " ); for( j = 0; j < 16; j++ ) log_error( "0x%02x ", g[ j ] ); return -1; } } return 0; } void bitselect_verify_fn( ExplicitType vecType, ExplicitType testVecType, unsigned int vecSize, void *inDataA, void *inDataB, void *inDataTest, void *outData ) { char *inA = (char *)inDataA, *inB = (char *)inDataB, *inT = (char *)inDataTest, *out = (char *)outData; size_t i, numBytes = get_explicit_type_size( vecType ); // Type is meaningless, this is all bitwise! for( i = 0; i < numBytes; i++ ) { out[ i ] = ( inA[ i ] & ~inT[ i ] ) | ( inB[ i ] & inT[ i ] ); } } int test_relational_bitselect(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kFloat, kDouble }; unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 }; unsigned int index, typeIndex; int retVal = 0; RandomSeed seed( gRandomSeed ); for( typeIndex = 0; typeIndex < 10; typeIndex++ ) { if ((vecType[typeIndex] == kLong || vecType[typeIndex] == kULong) && !gHasLong) continue; if (vecType[typeIndex] == kDouble) { if(!is_extension_available(device, "cl_khr_fp64")) { log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n"); continue; } else log_info("Testing doubles.\n"); } for( index = 0; vecSizes[ index ] != 0; index++ ) { // Test! if( test_select_kernel(context, queue, "bitselect", vecType[ typeIndex ], vecSizes[ index ], vecType[typeIndex], bitselect_verify_fn, seed ) != 0 ) { log_error( " Vector %s%d FAILED\n", get_explicit_type_name( vecType[ typeIndex ] ), vecSizes[ index ] ); retVal = -1; } } } return retVal; } void select_signed_verify_fn( ExplicitType vecType, ExplicitType testVecType, unsigned int vecSize, void *inDataA, void *inDataB, void *inDataTest, void *outData ) { bool yep = false; if (vecSize == 1) { switch( testVecType ) { case kChar: yep = *( (char *)inDataTest ) ? true : false; break; case kShort: yep = *( (short *)inDataTest ) ? true : false; break; case kInt: yep = *( (int *)inDataTest ) ? true : false; break; case kLong: yep = *( (cl_long *)inDataTest ) ? true : false; break; default: // Should never get here return; } } else { switch( testVecType ) { case kChar: yep = *( (char *)inDataTest ) & 0x80 ? true : false; break; case kShort: yep = *( (short *)inDataTest ) & 0x8000 ? true : false; break; case kInt: yep = *( (int *)inDataTest ) & 0x80000000L ? true : false; break; case kLong: yep = *( (cl_long *)inDataTest ) & 0x8000000000000000LL ? true : false; break; default: // Should never get here return; } } memcpy( outData, ( yep ) ? inDataB : inDataA, get_explicit_type_size( vecType ) ); } int test_relational_select_signed(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kFloat, kDouble }; ExplicitType testVecType[] = { kChar, kShort, kInt, kLong, kNumExplicitTypes }; unsigned int vecSizes[] = { 1, 2, 4, 8, 16, 0 }; unsigned int index, typeIndex, testTypeIndex; int retVal = 0; RandomSeed seed( gRandomSeed ); for( typeIndex = 0; typeIndex < 10; typeIndex++ ) { if ((vecType[typeIndex] == kLong || vecType[typeIndex] == kULong) && !gHasLong) continue; if (vecType[typeIndex] == kDouble) { if(!is_extension_available(device, "cl_khr_fp64")) { log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n"); continue; } else { log_info("Testing doubles.\n"); } } for( testTypeIndex = 0; testVecType[ testTypeIndex ] != kNumExplicitTypes; testTypeIndex++ ) { if( testVecType[ testTypeIndex ] != vecType[ typeIndex ] ) continue; for( index = 0; vecSizes[ index ] != 0; index++ ) { // Test! if( test_select_kernel(context, queue, "select", vecType[ typeIndex ], vecSizes[ index ], testVecType[ testTypeIndex ], select_signed_verify_fn, seed ) != 0 ) { log_error( " Vector %s%d, test vector %s%d FAILED\n", get_explicit_type_name( vecType[ typeIndex ] ), vecSizes[ index ], get_explicit_type_name( testVecType[ testTypeIndex ] ), vecSizes[ index ] ); retVal = -1; } } } } return retVal; } void select_unsigned_verify_fn( ExplicitType vecType, ExplicitType testVecType, unsigned int vecSize, void *inDataA, void *inDataB, void *inDataTest, void *outData ) { bool yep = false; if (vecSize == 1) { switch( testVecType ) { case kUChar: yep = *( (unsigned char *)inDataTest ) ? true : false; break; case kUShort: yep = *( (unsigned short *)inDataTest ) ? true : false; break; case kUInt: yep = *( (unsigned int *)inDataTest ) ? true : false; break; case kULong: yep = *( (cl_ulong *)inDataTest ) ? true : false; break; default: // Should never get here return; } } else { switch( testVecType ) { case kUChar: yep = *( (unsigned char *)inDataTest ) & 0x80 ? true : false; break; case kUShort: yep = *( (unsigned short *)inDataTest ) & 0x8000 ? true : false; break; case kUInt: yep = *( (unsigned int *)inDataTest ) & 0x80000000L ? true : false; break; case kULong: yep = *( (cl_ulong *)inDataTest ) & 0x8000000000000000LL ? true : false; break; default: // Should never get here return; } } memcpy( outData, ( yep ) ? inDataB : inDataA, get_explicit_type_size( vecType ) ); } int test_relational_select_unsigned(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kFloat, kDouble }; ExplicitType testVecType[] = { kUChar, kUShort, kUInt, kULong, kNumExplicitTypes }; unsigned int vecSizes[] = { 1, 2, 4, 8, 16, 0 }; unsigned int index, typeIndex, testTypeIndex; int retVal = 0; RandomSeed seed(gRandomSeed); for( typeIndex = 0; typeIndex < 10; typeIndex++ ) { if ((vecType[typeIndex] == kLong || vecType[typeIndex] == kULong) && !gHasLong) continue; if (vecType[typeIndex] == kDouble) { if(!is_extension_available(device, "cl_khr_fp64")) { log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n"); continue; } else { log_info("Testing doubles.\n"); } } for( testTypeIndex = 0; testVecType[ testTypeIndex ] != kNumExplicitTypes; testTypeIndex++ ) { if( testVecType[ testTypeIndex ] != vecType[ typeIndex ] ) continue; for( index = 0; vecSizes[ index ] != 0; index++ ) { // Test! if( test_select_kernel(context, queue, "select", vecType[ typeIndex ], vecSizes[ index ], testVecType[ testTypeIndex ], select_unsigned_verify_fn, seed ) != 0 ) { log_error( " Vector %s%d, test vector %s%d FAILED\n", get_explicit_type_name( vecType[ typeIndex ] ), vecSizes[ index ], get_explicit_type_name( testVecType[ testTypeIndex ] ), vecSizes[ index ] ); retVal = -1; } } } } return retVal; } extern int test_relational_isequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); extern int test_relational_isnotequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); extern int test_relational_isgreater_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); extern int test_relational_isgreaterequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); extern int test_relational_isless_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); extern int test_relational_islessequal_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); extern int test_relational_islessgreater_float(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); extern int test_relational_isequal_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); extern int test_relational_isnotequal_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); extern int test_relational_isgreater_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); extern int test_relational_isgreaterequal_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); extern int test_relational_isless_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); extern int test_relational_islessequal_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); extern int test_relational_islessgreater_double(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ); int test_relational_isequal(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { int err = 0; err |= test_relational_isequal_float( device, context, queue, numElements ); err |= test_relational_isequal_double( device, context, queue, numElements ); return err; } int test_relational_isnotequal(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { int err = 0; err |= test_relational_isnotequal_float( device, context, queue, numElements ); err |= test_relational_isnotequal_double( device, context, queue, numElements ); return err; } int test_relational_isgreater(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { int err = 0; err |= test_relational_isgreater_float( device, context, queue, numElements ); err |= test_relational_isgreater_double( device, context, queue, numElements ); return err; } int test_relational_isgreaterequal(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { int err = 0; err |= test_relational_isgreaterequal_float( device, context, queue, numElements ); err |= test_relational_isgreaterequal_double( device, context, queue, numElements ); return err; } int test_relational_isless(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { int err = 0; err |= test_relational_isless_float( device, context, queue, numElements ); err |= test_relational_isless_double( device, context, queue, numElements ); return err; } int test_relational_islessequal(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { int err = 0; err |= test_relational_islessequal_float( device, context, queue, numElements ); err |= test_relational_islessequal_double( device, context, queue, numElements ); return err; } int test_relational_islessgreater(cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { int err = 0; err |= test_relational_islessgreater_float( device, context, queue, numElements ); err |= test_relational_islessgreater_double( device, context, queue, numElements ); return err; }