1 //
2 // Copyright (c) 2017 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 // http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16 #include "testBase.h"
17 #include "harness/conversions.h"
18 #include "harness/typeWrappers.h"
19
20 #if !defined (__APPLE__)
21 #include <CL/cl_gl.h>
22 #endif
23
24 static const char *bufferKernelPattern =
25 "__kernel void sample_test( __global %s%s *source, __global %s%s *clDest, __global %s%s *glDest )\n"
26 "{\n"
27 " int tid = get_global_id(0);\n"
28 " clDest[ tid ] = source[ tid ] + (%s%s)(1);\n"
29 " glDest[ tid ] = source[ tid ] + (%s%s)(2);\n"
30 "}\n";
31
32 #define TYPE_CASE( enum, type, range, offset ) \
33 case enum: \
34 { \
35 cl_##type *ptr = (cl_##type *)outData; \
36 for( i = 0; i < count; i++ ) \
37 ptr[ i ] = (cl_##type)( ( genrand_int32(d) & range ) - offset ); \
38 break; \
39 }
40
gen_input_data(ExplicitType type,size_t count,MTdata d,void * outData)41 void gen_input_data( ExplicitType type, size_t count, MTdata d, void *outData )
42 {
43 size_t i;
44
45 switch( type )
46 {
47 case kBool:
48 {
49 bool *boolPtr = (bool *)outData;
50 for( i = 0; i < count; i++ )
51 {
52 boolPtr[i] = ( genrand_int32(d) & 1 ) ? true : false;
53 }
54 break;
55 }
56
57 TYPE_CASE( kChar, char, 250, 127 )
58 TYPE_CASE( kUChar, uchar, 250, 0 )
59 TYPE_CASE( kShort, short, 65530, 32767 )
60 TYPE_CASE( kUShort, ushort, 65530, 0 )
61 TYPE_CASE( kInt, int, 0x0fffffff, 0x70000000 )
62 TYPE_CASE( kUInt, uint, 0x0fffffff, 0 )
63
64 case kLong:
65 {
66 cl_long *longPtr = (cl_long *)outData;
67 for( i = 0; i < count; i++ )
68 {
69 longPtr[i] = (cl_long)genrand_int32(d) | ( (cl_ulong)genrand_int32(d) << 32 );
70 }
71 break;
72 }
73
74 case kULong:
75 {
76 cl_ulong *ulongPtr = (cl_ulong *)outData;
77 for( i = 0; i < count; i++ )
78 {
79 ulongPtr[i] = (cl_ulong)genrand_int32(d) | ( (cl_ulong)genrand_int32(d) << 32 );
80 }
81 break;
82 }
83
84 case kFloat:
85 {
86 cl_float *floatPtr = (float *)outData;
87 for( i = 0; i < count; i++ )
88 floatPtr[i] = get_random_float( -100000.f, 100000.f, d );
89 break;
90 }
91
92 default:
93 log_error( "ERROR: Invalid type passed in to generate_random_data!\n" );
94 break;
95 }
96 }
97
98 #define INC_CASE( enum, type ) \
99 case enum: \
100 { \
101 cl_##type *src = (cl_##type *)inData; \
102 cl_##type *dst = (cl_##type *)outData; \
103 *dst = *src + 1; \
104 break; \
105 }
106
get_incremented_value(void * inData,void * outData,ExplicitType type)107 void get_incremented_value( void *inData, void *outData, ExplicitType type )
108 {
109 switch( type )
110 {
111 INC_CASE( kChar, char )
112 INC_CASE( kUChar, uchar )
113 INC_CASE( kShort, short )
114 INC_CASE( kUShort, ushort )
115 INC_CASE( kInt, int )
116 INC_CASE( kUInt, uint )
117 INC_CASE( kLong, long )
118 INC_CASE( kULong, ulong )
119 INC_CASE( kFloat, float )
120 default:
121 break;
122 }
123 }
124
test_buffer_kernel(cl_context context,cl_command_queue queue,ExplicitType vecType,size_t vecSize,int numElements,int validate_only,MTdata d)125 int test_buffer_kernel(cl_context context, cl_command_queue queue, ExplicitType vecType, size_t vecSize, int numElements, int validate_only, MTdata d)
126 {
127 clProgramWrapper program;
128 clKernelWrapper kernel;
129 clMemWrapper streams[ 3 ];
130 size_t dataSize = numElements * 16 * sizeof(cl_long);
131 #if !(defined(_WIN32) && defined(_MSC_VER))
132 cl_long inData[numElements * 16], outDataCL[numElements * 16], outDataGL[ numElements * 16 ];
133 #else
134 cl_long* inData = (cl_long*)_malloca(dataSize);
135 cl_long* outDataCL = (cl_long*)_malloca(dataSize);
136 cl_long* outDataGL = (cl_long*)_malloca(dataSize);
137 #endif
138 glBufferWrapper inGLBuffer, outGLBuffer;
139 int i;
140 size_t bufferSize;
141
142 int error;
143 size_t threads[1], localThreads[1];
144 char kernelSource[10240];
145 char *programPtr;
146 char sizeName[4];
147
148 /* Create the source */
149 if( vecSize == 1 )
150 sizeName[ 0 ] = 0;
151 else
152 sprintf( sizeName, "%d", (int)vecSize );
153
154 sprintf( kernelSource, bufferKernelPattern, get_explicit_type_name( vecType ), sizeName,
155 get_explicit_type_name( vecType ), sizeName,
156 get_explicit_type_name( vecType ), sizeName,
157 get_explicit_type_name( vecType ), sizeName,
158 get_explicit_type_name( vecType ), sizeName );
159
160 /* Create kernels */
161 programPtr = kernelSource;
162 if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "sample_test" ) )
163 {
164 return -1;
165 }
166
167 bufferSize = numElements * vecSize * get_explicit_type_size( vecType );
168
169 /* Generate some almost-random input data */
170 gen_input_data( vecType, vecSize * numElements, d, inData );
171 memset( outDataCL, 0, dataSize );
172 memset( outDataGL, 0, dataSize );
173
174 /* Generate some GL buffers to go against */
175 glGenBuffers( 1, &inGLBuffer );
176 glGenBuffers( 1, &outGLBuffer );
177
178 glBindBuffer( GL_ARRAY_BUFFER, inGLBuffer );
179 glBufferData( GL_ARRAY_BUFFER, bufferSize, inData, GL_STATIC_DRAW );
180
181 // Note: we need to bind the output buffer, even though we don't care about its values yet,
182 // because CL needs it to get the buffer size
183 glBindBuffer( GL_ARRAY_BUFFER, outGLBuffer );
184 glBufferData( GL_ARRAY_BUFFER, bufferSize, outDataGL, GL_STATIC_DRAW );
185
186 glBindBuffer( GL_ARRAY_BUFFER, 0 );
187 glFlush();
188
189
190 /* Generate some streams. The first and last ones are GL, middle one just vanilla CL */
191 streams[ 0 ] = (*clCreateFromGLBuffer_ptr)( context, CL_MEM_READ_ONLY, inGLBuffer, &error );
192 test_error( error, "Unable to create input GL buffer" );
193
194 streams[ 1 ] = clCreateBuffer( context, CL_MEM_READ_WRITE, bufferSize, NULL, &error );
195 test_error( error, "Unable to create output CL buffer" );
196
197 streams[ 2 ] = (*clCreateFromGLBuffer_ptr)( context, CL_MEM_WRITE_ONLY, outGLBuffer, &error );
198 test_error( error, "Unable to create output GL buffer" );
199
200
201 /* Validate the info */
202 if (validate_only) {
203 int result = (CheckGLObjectInfo(streams[0], CL_GL_OBJECT_BUFFER, (GLuint)inGLBuffer, (GLenum)0, 0) |
204 CheckGLObjectInfo(streams[2], CL_GL_OBJECT_BUFFER, (GLuint)outGLBuffer, (GLenum)0, 0) );
205 for(i=0;i<3;i++)
206 {
207 clReleaseMemObject(streams[i]);
208 streams[i] = NULL;
209 }
210
211 glDeleteBuffers(1, &inGLBuffer); inGLBuffer = 0;
212 glDeleteBuffers(1, &outGLBuffer); outGLBuffer = 0;
213
214 return result;
215 }
216
217 /* Assign streams and execute */
218 for( int i = 0; i < 3; i++ )
219 {
220 error = clSetKernelArg( kernel, i, sizeof( streams[ i ] ), &streams[ i ] );
221 test_error( error, "Unable to set kernel arguments" );
222 }
223 error = (*clEnqueueAcquireGLObjects_ptr)( queue, 1, &streams[ 0 ], 0, NULL, NULL);
224 test_error( error, "Unable to acquire GL obejcts");
225 error = (*clEnqueueAcquireGLObjects_ptr)( queue, 1, &streams[ 2 ], 0, NULL, NULL);
226 test_error( error, "Unable to acquire GL obejcts");
227
228 /* Run the kernel */
229 threads[0] = numElements;
230
231 error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
232 test_error( error, "Unable to get work group size to use" );
233
234 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
235 test_error( error, "Unable to execute test kernel" );
236
237 error = (*clEnqueueReleaseGLObjects_ptr)( queue, 1, &streams[ 0 ], 0, NULL, NULL );
238 test_error(error, "clEnqueueReleaseGLObjects failed");
239 error = (*clEnqueueReleaseGLObjects_ptr)( queue, 1, &streams[ 2 ], 0, NULL, NULL );
240 test_error(error, "clEnqueueReleaseGLObjects failed");
241
242 // Get the results from both CL and GL and make sure everything looks correct
243 error = clEnqueueReadBuffer( queue, streams[ 1 ], CL_TRUE, 0, bufferSize, outDataCL, 0, NULL, NULL );
244 test_error( error, "Unable to read output CL array!" );
245
246 glBindBuffer( GL_ARRAY_BUFFER, outGLBuffer );
247 void *glMem = glMapBuffer( GL_ARRAY_BUFFER, GL_READ_ONLY );
248 memcpy( outDataGL, glMem, bufferSize );
249 glUnmapBuffer( GL_ARRAY_BUFFER );
250
251 char *inP = (char *)inData, *glP = (char *)outDataGL, *clP = (char *)outDataCL;
252 error = 0;
253 for( size_t i = 0; i < numElements * vecSize; i++ )
254 {
255 cl_long expectedCLValue, expectedGLValue;
256 get_incremented_value( inP, &expectedCLValue, vecType );
257 get_incremented_value( &expectedCLValue, &expectedGLValue, vecType );
258
259 if( memcmp( clP, &expectedCLValue, get_explicit_type_size( vecType ) ) != 0 )
260 {
261 char scratch[ 64 ];
262 log_error( "ERROR: Data sample %d from the CL output did not validate!\n", (int)i );
263 log_error( "\t Input: %s\n", GetDataVectorString( inP, get_explicit_type_size( vecType ), 1, scratch ) );
264 log_error( "\tExpected: %s\n", GetDataVectorString( &expectedCLValue, get_explicit_type_size( vecType ), 1, scratch ) );
265 log_error( "\t Actual: %s\n", GetDataVectorString( clP, get_explicit_type_size( vecType ), 1, scratch ) );
266 error = -1;
267 }
268
269 if( memcmp( glP, &expectedGLValue, get_explicit_type_size( vecType ) ) != 0 )
270 {
271 char scratch[ 64 ];
272 log_error( "ERROR: Data sample %d from the GL output did not validate!\n", (int)i );
273 log_error( "\t Input: %s\n", GetDataVectorString( inP, get_explicit_type_size( vecType ), 1, scratch ) );
274 log_error( "\tExpected: %s\n", GetDataVectorString( &expectedGLValue, get_explicit_type_size( vecType ), 1, scratch ) );
275 log_error( "\t Actual: %s\n", GetDataVectorString( glP, get_explicit_type_size( vecType ), 1, scratch ) );
276 error = -1;
277 }
278
279 if( error )
280 return error;
281
282 inP += get_explicit_type_size( vecType );
283 glP += get_explicit_type_size( vecType );
284 clP += get_explicit_type_size( vecType );
285 }
286
287 for(i=0;i<3;i++)
288 {
289 clReleaseMemObject(streams[i]);
290 streams[i] = NULL;
291 }
292
293 glDeleteBuffers(1, &inGLBuffer); inGLBuffer = 0;
294 glDeleteBuffers(1, &outGLBuffer); outGLBuffer = 0;
295
296 return 0;
297 }
298
test_buffers(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)299 int test_buffers( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
300 {
301 ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kFloat, kNumExplicitTypes };
302 unsigned int vecSizes[] = { 1, 2, 4, 8, 16, 0 };
303 unsigned int index, typeIndex;
304 int retVal = 0;
305 RandomSeed seed(gRandomSeed);
306
307
308 for( typeIndex = 0; vecType[ typeIndex ] != kNumExplicitTypes; typeIndex++ )
309 {
310 for( index = 0; vecSizes[ index ] != 0; index++ )
311 {
312 // Test!
313 if( test_buffer_kernel( context, queue, vecType[ typeIndex ], vecSizes[ index ], numElements, 0, seed) != 0 )
314 {
315 char sizeNames[][ 4 ] = { "", "", "2", "", "4", "", "", "", "8", "", "", "", "", "", "", "", "16" };
316 log_error( " Buffer test %s%s FAILED\n", get_explicit_type_name( vecType[ typeIndex ] ), sizeNames[ vecSizes[ index ] ] );
317 retVal++;
318 }
319 }
320 }
321
322 return retVal;
323
324 }
325
326
test_buffers_getinfo(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)327 int test_buffers_getinfo( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
328 {
329 ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kFloat, kNumExplicitTypes };
330 unsigned int vecSizes[] = { 1, 2, 4, 8, 16, 0 };
331 unsigned int index, typeIndex;
332 int retVal = 0;
333 RandomSeed seed( gRandomSeed );
334
335
336 for( typeIndex = 0; vecType[ typeIndex ] != kNumExplicitTypes; typeIndex++ )
337 {
338 for( index = 0; vecSizes[ index ] != 0; index++ )
339 {
340 // Test!
341 if( test_buffer_kernel( context, queue, vecType[ typeIndex ], vecSizes[ index ], numElements, 1, seed ) != 0 )
342 {
343 char sizeNames[][ 4 ] = { "", "", "2", "", "4", "", "", "", "8", "", "", "", "", "", "", "", "16" };
344 log_error( " Buffer test %s%s FAILED\n", get_explicit_type_name( vecType[ typeIndex ] ), sizeNames[ vecSizes[ index ] ] );
345 retVal++;
346 }
347 }
348 }
349
350 return retVal;
351
352 }
353
354
355
356