// // 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" #define TEST_SIZE 512 enum OpKonstants { kIncrement = 0, kDecrement, kBoth }; const char *testKernel = "__kernel void test( __global %s *inOut, __global char * control )\n" "{\n" " size_t tid = get_global_id(0);\n" "\n" " %s%s inOutVal = %s;\n" "\n" " if( control[tid] == 0 )\n" " inOutVal++;\n" " else if( control[tid] == 1 )\n" " ++inOutVal;\n" " else if( control[tid] == 2 )\n" " inOutVal--;\n" " else // if( control[tid] == 3 )\n" " --inOutVal;\n" "\n" " %s;\n" "}\n"; typedef int (*OpVerifyFn)( void * actualPtr, void * inputPtr, size_t vecSize, size_t numVecs, cl_char * controls ); int test_unary_op( cl_command_queue queue, cl_context context, OpKonstants whichOp, ExplicitType vecType, size_t vecSize, MTdata d, OpVerifyFn verifyFn ) { clProgramWrapper program; clKernelWrapper kernel; clMemWrapper streams[2]; cl_long inData[TEST_SIZE * 16], outData[TEST_SIZE * 16]; cl_char controlData[TEST_SIZE]; int error; size_t i; size_t threads[1], localThreads[1]; char kernelSource[10240]; char *programPtr; // Create the source char loadLine[ 1024 ], storeLine[ 1024 ]; if( vecSize == 1 ) { sprintf( loadLine, "inOut[tid]" ); sprintf( storeLine, "inOut[tid] = inOutVal" ); } else { sprintf( loadLine, "vload%ld( tid, inOut )", vecSize ); sprintf( storeLine, "vstore%ld( inOutVal, tid, inOut )", vecSize ); } char sizeNames[][4] = { "", "", "2", "3", "4", "", "", "", "8", "", "", "", "", "", "", "", "16" }; sprintf( kernelSource, testKernel, get_explicit_type_name( vecType ), /*sizeNames[ vecSize ],*/ get_explicit_type_name( vecType ), sizeNames[ vecSize ], loadLine, storeLine ); // Create the kernel programPtr = kernelSource; if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "test" ) ) { log_error( "ERROR: Unable to create test program!\n" ); return -1; } // Generate two streams. The first is our random data to test against, the second is our control stream generate_random_data( vecType, vecSize * TEST_SIZE, d, inData ); streams[0] = clCreateBuffer( context, CL_MEM_COPY_HOST_PTR, get_explicit_type_size(vecType) * vecSize * TEST_SIZE, inData, &error); test_error( error, "Creating input data array failed" ); cl_uint bits; for( i = 0; i < TEST_SIZE; i++ ) { size_t which = i & 7; if( which == 0 ) bits = genrand_int32(d); controlData[ i ] = ( bits >> ( which << 1 ) ) & 0x03; if( whichOp == kDecrement ) // For sub ops, the min control value is 2. Otherwise, it's 0 controlData[ i ] |= 0x02; else if( whichOp == kIncrement ) // For addition ops, the max control value is 1. Otherwise, it's 3 controlData[ i ] &= ~0x02; } streams[1] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(controlData), controlData, &error); test_error( error, "Unable to create control stream" ); // 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" ); // Read the results error = clEnqueueReadBuffer( queue, streams[0], CL_TRUE, 0, get_explicit_type_size( vecType ) * TEST_SIZE * vecSize, outData, 0, NULL, NULL ); test_error( error, "Unable to read output array!" ); // Now verify the results return verifyFn( outData, inData, vecSize, TEST_SIZE, controlData ); } template int VerifyFn( void * actualPtr, void * inputPtr, size_t vecSize, size_t numVecs, cl_char * controls ) { T * actualData = (T *)actualPtr; T * inputData = (T *)inputPtr; size_t index = 0; for( size_t i = 0; i < numVecs; i++ ) { for( size_t j = 0; j < vecSize; j++, index++ ) { T nextVal = inputData[ index ]; if( controls[ i ] & 0x02 ) nextVal--; else nextVal++; if( actualData[ index ] != nextVal ) { log_error( "ERROR: Validation failed on vector %ld:%ld (expected %lld, got %lld)", i, j, (cl_long)nextVal, (cl_long)actualData[ index ] ); return -1; } } } return 0; } int test_unary_op_set( cl_command_queue queue, cl_context context, OpKonstants whichOp ) { ExplicitType types[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kNumExplicitTypes }; OpVerifyFn verifys[] = { VerifyFn, VerifyFn, VerifyFn, VerifyFn, VerifyFn, VerifyFn, VerifyFn, VerifyFn, NULL }; unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 }; unsigned int index, typeIndex; int retVal = 0; RandomSeed seed(gRandomSeed ); for( typeIndex = 0; types[ typeIndex ] != kNumExplicitTypes; typeIndex++ ) { if ((types[ typeIndex ] == kLong || types[ typeIndex ] == kULong) && !gHasLong) continue; for( index = 0; vecSizes[ index ] != 0; index++ ) { if( test_unary_op( queue, context, whichOp, types[ typeIndex ], vecSizes[ index ], seed, verifys[ typeIndex ] ) != 0 ) { log_error( " Vector %s%d FAILED\n", get_explicit_type_name( types[ typeIndex ] ), vecSizes[ index ] ); retVal = -1; } } } return retVal; } int test_unary_ops_full(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { return test_unary_op_set( queue, context, kBoth ); } int test_unary_ops_increment(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { return test_unary_op_set( queue, context, kIncrement ); } int test_unary_ops_decrement(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { return test_unary_op_set( queue, context, kDecrement ); }