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