| // |
| // 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<typename T> 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<cl_char>, VerifyFn<cl_uchar>, VerifyFn<cl_short>, VerifyFn<cl_ushort>, VerifyFn<cl_int>, VerifyFn<cl_uint>, VerifyFn<cl_long>, VerifyFn<cl_ulong>, 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 ); |
| } |