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