// // 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" #if !defined (__APPLE__) #include #endif static const char *bufferKernelPattern = "__kernel void sample_test( __global %s%s *source, __global %s%s *clDest, __global %s%s *glDest )\n" "{\n" " int tid = get_global_id(0);\n" " clDest[ tid ] = source[ tid ] + (%s%s)(1);\n" " glDest[ tid ] = source[ tid ] + (%s%s)(2);\n" "}\n"; #define TYPE_CASE( enum, type, range, offset ) \ case enum: \ { \ cl_##type *ptr = (cl_##type *)outData; \ for( i = 0; i < count; i++ ) \ ptr[ i ] = (cl_##type)( ( genrand_int32(d) & range ) - offset ); \ break; \ } void gen_input_data( ExplicitType type, size_t count, MTdata d, void *outData ) { size_t i; switch( type ) { case kBool: { bool *boolPtr = (bool *)outData; for( i = 0; i < count; i++ ) { boolPtr[i] = ( genrand_int32(d) & 1 ) ? true : false; } break; } TYPE_CASE( kChar, char, 250, 127 ) TYPE_CASE( kUChar, uchar, 250, 0 ) TYPE_CASE( kShort, short, 65530, 32767 ) TYPE_CASE( kUShort, ushort, 65530, 0 ) TYPE_CASE( kInt, int, 0x0fffffff, 0x70000000 ) TYPE_CASE( kUInt, uint, 0x0fffffff, 0 ) case kLong: { cl_long *longPtr = (cl_long *)outData; for( i = 0; i < count; i++ ) { longPtr[i] = (cl_long)genrand_int32(d) | ( (cl_ulong)genrand_int32(d) << 32 ); } break; } case kULong: { cl_ulong *ulongPtr = (cl_ulong *)outData; for( i = 0; i < count; i++ ) { ulongPtr[i] = (cl_ulong)genrand_int32(d) | ( (cl_ulong)genrand_int32(d) << 32 ); } break; } case kFloat: { cl_float *floatPtr = (float *)outData; for( i = 0; i < count; i++ ) floatPtr[i] = get_random_float( -100000.f, 100000.f, d ); break; } default: log_error( "ERROR: Invalid type passed in to generate_random_data!\n" ); break; } } #define INC_CASE( enum, type ) \ case enum: \ { \ cl_##type *src = (cl_##type *)inData; \ cl_##type *dst = (cl_##type *)outData; \ *dst = *src + 1; \ break; \ } void get_incremented_value( void *inData, void *outData, ExplicitType type ) { switch( type ) { INC_CASE( kChar, char ) INC_CASE( kUChar, uchar ) INC_CASE( kShort, short ) INC_CASE( kUShort, ushort ) INC_CASE( kInt, int ) INC_CASE( kUInt, uint ) INC_CASE( kLong, long ) INC_CASE( kULong, ulong ) INC_CASE( kFloat, float ) default: break; } } int test_buffer_kernel(cl_context context, cl_command_queue queue, ExplicitType vecType, size_t vecSize, int numElements, int validate_only, MTdata d) { clProgramWrapper program; clKernelWrapper kernel; clMemWrapper streams[ 3 ]; size_t dataSize = numElements * 16 * sizeof(cl_long); #if !(defined(_WIN32) && defined(_MSC_VER)) cl_long inData[numElements * 16], outDataCL[numElements * 16], outDataGL[ numElements * 16 ]; #else cl_long* inData = (cl_long*)_malloca(dataSize); cl_long* outDataCL = (cl_long*)_malloca(dataSize); cl_long* outDataGL = (cl_long*)_malloca(dataSize); #endif glBufferWrapper inGLBuffer, outGLBuffer; int i; size_t bufferSize; int error; size_t threads[1], localThreads[1]; char kernelSource[10240]; char *programPtr; char sizeName[4]; /* Create the source */ if( vecSize == 1 ) sizeName[ 0 ] = 0; else sprintf( sizeName, "%d", (int)vecSize ); sprintf( kernelSource, bufferKernelPattern, get_explicit_type_name( vecType ), sizeName, get_explicit_type_name( vecType ), sizeName, get_explicit_type_name( vecType ), sizeName, get_explicit_type_name( vecType ), sizeName, get_explicit_type_name( vecType ), sizeName ); /* Create kernels */ programPtr = kernelSource; if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "sample_test" ) ) { return -1; } bufferSize = numElements * vecSize * get_explicit_type_size( vecType ); /* Generate some almost-random input data */ gen_input_data( vecType, vecSize * numElements, d, inData ); memset( outDataCL, 0, dataSize ); memset( outDataGL, 0, dataSize ); /* Generate some GL buffers to go against */ glGenBuffers( 1, &inGLBuffer ); glGenBuffers( 1, &outGLBuffer ); glBindBuffer( GL_ARRAY_BUFFER, inGLBuffer ); glBufferData( GL_ARRAY_BUFFER, bufferSize, inData, GL_STATIC_DRAW ); // Note: we need to bind the output buffer, even though we don't care about its values yet, // because CL needs it to get the buffer size glBindBuffer( GL_ARRAY_BUFFER, outGLBuffer ); glBufferData( GL_ARRAY_BUFFER, bufferSize, outDataGL, GL_STATIC_DRAW ); glBindBuffer( GL_ARRAY_BUFFER, 0 ); glFinish(); /* Generate some streams. The first and last ones are GL, middle one just vanilla CL */ streams[ 0 ] = (*clCreateFromGLBuffer_ptr)( context, CL_MEM_READ_ONLY, inGLBuffer, &error ); test_error( error, "Unable to create input GL buffer" ); streams[ 1 ] = clCreateBuffer( context, CL_MEM_READ_WRITE, bufferSize, NULL, &error ); test_error( error, "Unable to create output CL buffer" ); streams[ 2 ] = (*clCreateFromGLBuffer_ptr)( context, CL_MEM_WRITE_ONLY, outGLBuffer, &error ); test_error( error, "Unable to create output GL buffer" ); /* Validate the info */ if (validate_only) { int result = (CheckGLObjectInfo(streams[0], CL_GL_OBJECT_BUFFER, (GLuint)inGLBuffer, (GLenum)0, 0) | CheckGLObjectInfo(streams[2], CL_GL_OBJECT_BUFFER, (GLuint)outGLBuffer, (GLenum)0, 0) ); for(i=0;i<3;i++) { clReleaseMemObject(streams[i]); streams[i] = NULL; } glDeleteBuffers(1, &inGLBuffer); inGLBuffer = 0; glDeleteBuffers(1, &outGLBuffer); outGLBuffer = 0; return result; } /* Assign streams and execute */ for( int i = 0; i < 3; i++ ) { error = clSetKernelArg( kernel, i, sizeof( streams[ i ] ), &streams[ i ] ); test_error( error, "Unable to set kernel arguments" ); } error = (*clEnqueueAcquireGLObjects_ptr)( queue, 1, &streams[ 0 ], 0, NULL, NULL); test_error( error, "Unable to acquire GL obejcts"); error = (*clEnqueueAcquireGLObjects_ptr)( queue, 1, &streams[ 2 ], 0, NULL, NULL); test_error( error, "Unable to acquire GL obejcts"); /* Run the kernel */ threads[0] = numElements; 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" ); error = (*clEnqueueReleaseGLObjects_ptr)( queue, 1, &streams[ 0 ], 0, NULL, NULL ); test_error(error, "clEnqueueReleaseGLObjects failed"); error = (*clEnqueueReleaseGLObjects_ptr)( queue, 1, &streams[ 2 ], 0, NULL, NULL ); test_error(error, "clEnqueueReleaseGLObjects failed"); // Get the results from both CL and GL and make sure everything looks correct error = clEnqueueReadBuffer( queue, streams[ 1 ], CL_TRUE, 0, bufferSize, outDataCL, 0, NULL, NULL ); test_error( error, "Unable to read output CL array!" ); glBindBuffer( GL_ARRAY_BUFFER, outGLBuffer ); void *glMem = glMapBuffer( GL_ARRAY_BUFFER, GL_READ_ONLY ); memcpy( outDataGL, glMem, bufferSize ); glUnmapBuffer( GL_ARRAY_BUFFER ); char *inP = (char *)inData, *glP = (char *)outDataGL, *clP = (char *)outDataCL; error = 0; for( size_t i = 0; i < numElements * vecSize; i++ ) { cl_long expectedCLValue, expectedGLValue; get_incremented_value( inP, &expectedCLValue, vecType ); get_incremented_value( &expectedCLValue, &expectedGLValue, vecType ); if( memcmp( clP, &expectedCLValue, get_explicit_type_size( vecType ) ) != 0 ) { char scratch[ 64 ]; log_error( "ERROR: Data sample %d from the CL output did not validate!\n", (int)i ); log_error( "\t Input: %s\n", GetDataVectorString( inP, get_explicit_type_size( vecType ), 1, scratch ) ); log_error( "\tExpected: %s\n", GetDataVectorString( &expectedCLValue, get_explicit_type_size( vecType ), 1, scratch ) ); log_error( "\t Actual: %s\n", GetDataVectorString( clP, get_explicit_type_size( vecType ), 1, scratch ) ); error = -1; } if( memcmp( glP, &expectedGLValue, get_explicit_type_size( vecType ) ) != 0 ) { char scratch[ 64 ]; log_error( "ERROR: Data sample %d from the GL output did not validate!\n", (int)i ); log_error( "\t Input: %s\n", GetDataVectorString( inP, get_explicit_type_size( vecType ), 1, scratch ) ); log_error( "\tExpected: %s\n", GetDataVectorString( &expectedGLValue, get_explicit_type_size( vecType ), 1, scratch ) ); log_error( "\t Actual: %s\n", GetDataVectorString( glP, get_explicit_type_size( vecType ), 1, scratch ) ); error = -1; } if( error ) return error; inP += get_explicit_type_size( vecType ); glP += get_explicit_type_size( vecType ); clP += get_explicit_type_size( vecType ); } for(i=0;i<3;i++) { clReleaseMemObject(streams[i]); streams[i] = NULL; } glDeleteBuffers(1, &inGLBuffer); inGLBuffer = 0; glDeleteBuffers(1, &outGLBuffer); outGLBuffer = 0; return 0; } int test_buffers( cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kFloat, kNumExplicitTypes }; unsigned int vecSizes[] = { 1, 2, 4, 8, 16, 0 }; unsigned int index, typeIndex; int retVal = 0; RandomSeed seed(gRandomSeed); for( typeIndex = 0; vecType[ typeIndex ] != kNumExplicitTypes; typeIndex++ ) { for( index = 0; vecSizes[ index ] != 0; index++ ) { // Test! if( test_buffer_kernel( context, queue, vecType[ typeIndex ], vecSizes[ index ], numElements, 0, seed) != 0 ) { char sizeNames[][ 4 ] = { "", "", "2", "", "4", "", "", "", "8", "", "", "", "", "", "", "", "16" }; log_error( " Buffer test %s%s FAILED\n", get_explicit_type_name( vecType[ typeIndex ] ), sizeNames[ vecSizes[ index ] ] ); retVal++; } } } return retVal; } int test_buffers_getinfo( cl_device_id device, cl_context context, cl_command_queue queue, int numElements ) { ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kFloat, kNumExplicitTypes }; unsigned int vecSizes[] = { 1, 2, 4, 8, 16, 0 }; unsigned int index, typeIndex; int retVal = 0; RandomSeed seed( gRandomSeed ); for( typeIndex = 0; vecType[ typeIndex ] != kNumExplicitTypes; typeIndex++ ) { for( index = 0; vecSizes[ index ] != 0; index++ ) { // Test! if( test_buffer_kernel( context, queue, vecType[ typeIndex ], vecSizes[ index ], numElements, 1, seed ) != 0 ) { char sizeNames[][ 4 ] = { "", "", "2", "", "4", "", "", "", "8", "", "", "", "", "", "", "", "16" }; log_error( " Buffer test %s%s FAILED\n", get_explicit_type_name( vecType[ typeIndex ] ), sizeNames[ vecSizes[ index ] ] ); retVal++; } } } return retVal; }