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/typeWrappers.h"
18 #include "harness/conversions.h"
19
20 const char *sample_single_test_kernel[] = {
21 "__kernel void sample_test(__global float *src, __global int *dst)\n"
22 "{\n"
23 " int tid = get_global_id(0);\n"
24 "\n"
25 " dst[tid] = (int)src[tid];\n"
26 "\n"
27 "}\n" };
28
29 const char *sample_struct_array_test_kernel[] = {
30 "typedef struct {\n"
31 "int A;\n"
32 "int B;\n"
33 "} input_pair_t;\n"
34 "\n"
35 "__kernel void sample_test(__global input_pair_t *src, __global int *dst)\n"
36 "{\n"
37 " int tid = get_global_id(0);\n"
38 "\n"
39 " dst[tid] = src[tid].A + src[tid].B;\n"
40 "\n"
41 "}\n" };
42
43 const char *sample_const_test_kernel[] = {
44 "__kernel void sample_test(__constant int *src1, __constant int *src2, __global int *dst)\n"
45 "{\n"
46 " int tid = get_global_id(0);\n"
47 "\n"
48 " dst[tid] = src1[tid] + src2[tid];\n"
49 "\n"
50 "}\n" };
51
52 const char *sample_const_global_test_kernel[] = {
53 "__constant int addFactor = 1024;\n"
54 "__kernel void sample_test(__global int *src1, __global int *dst)\n"
55 "{\n"
56 " int tid = get_global_id(0);\n"
57 "\n"
58 " dst[tid] = src1[tid] + addFactor;\n"
59 "\n"
60 "}\n" };
61
62 const char *sample_two_kernel_program[] = {
63 "__kernel void sample_test(__global float *src, __global int *dst)\n"
64 "{\n"
65 " int tid = get_global_id(0);\n"
66 "\n"
67 " dst[tid] = (int)src[tid];\n"
68 "\n"
69 "}\n",
70 "__kernel void sample_test2(__global int *src, __global float *dst)\n"
71 "{\n"
72 " int tid = get_global_id(0);\n"
73 "\n"
74 " dst[tid] = (float)src[tid];\n"
75 "\n"
76 "}\n" };
77
78
79
80
test_get_kernel_info(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)81 int test_get_kernel_info(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
82 {
83 int error;
84 cl_program program, testProgram;
85 cl_context testContext;
86 cl_kernel kernel;
87 cl_char name[ 512 ];
88 cl_uint numArgs, numInstances;
89 size_t paramSize;
90
91
92 /* Create reference */
93 if( create_single_kernel_helper( context, &program, &kernel, 1, sample_single_test_kernel, "sample_test" ) != 0 )
94 {
95 return -1;
96 }
97
98 error = clGetKernelInfo( kernel, CL_KERNEL_FUNCTION_NAME, 0, NULL, ¶mSize );
99 test_error( error, "Unable to get kernel function name param size" );
100 if( paramSize != strlen( "sample_test" ) + 1 )
101 {
102 log_error( "ERROR: Kernel function name param returns invalid size (expected %d, got %d)\n", (int)strlen( "sample_test" ) + 1, (int)paramSize );
103 return -1;
104 }
105
106 error = clGetKernelInfo( kernel, CL_KERNEL_FUNCTION_NAME, sizeof( name ), name, NULL );
107 test_error( error, "Unable to get kernel function name" );
108 if( strcmp( (char *)name, "sample_test" ) != 0 )
109 {
110 log_error( "ERROR: Kernel function name returned invalid value (expected sample_test, got %s)\n", (char *)name );
111 return -1;
112 }
113
114
115 error = clGetKernelInfo( kernel, CL_KERNEL_NUM_ARGS, 0, NULL, ¶mSize );
116 test_error( error, "Unable to get kernel arg count param size" );
117 if( paramSize != sizeof( numArgs ) )
118 {
119 log_error( "ERROR: Kernel arg count param returns invalid size (expected %d, got %d)\n", (int)sizeof( numArgs ), (int)paramSize );
120 return -1;
121 }
122
123 error = clGetKernelInfo( kernel, CL_KERNEL_NUM_ARGS, sizeof( numArgs ), &numArgs, NULL );
124 test_error( error, "Unable to get kernel arg count" );
125 if( numArgs != 2 )
126 {
127 log_error( "ERROR: Kernel arg count returned invalid value (expected %d, got %d)\n", 2, numArgs );
128 return -1;
129 }
130
131
132 error = clGetKernelInfo( kernel, CL_KERNEL_REFERENCE_COUNT, 0, NULL, ¶mSize );
133 test_error( error, "Unable to get kernel reference count param size" );
134 if( paramSize != sizeof( numInstances ) )
135 {
136 log_error( "ERROR: Kernel reference count param returns invalid size (expected %d, got %d)\n", (int)sizeof( numInstances ), (int)paramSize );
137 return -1;
138 }
139
140 error = clGetKernelInfo( kernel, CL_KERNEL_REFERENCE_COUNT, sizeof( numInstances ), &numInstances, NULL );
141 test_error( error, "Unable to get kernel reference count" );
142
143
144 error = clGetKernelInfo( kernel, CL_KERNEL_PROGRAM, 0, NULL, ¶mSize );
145 test_error( error, "Unable to get kernel program param size" );
146 if( paramSize != sizeof( testProgram ) )
147 {
148 log_error( "ERROR: Kernel program param returns invalid size (expected %d, got %d)\n", (int)sizeof( testProgram ), (int)paramSize );
149 return -1;
150 }
151
152 error = clGetKernelInfo( kernel, CL_KERNEL_PROGRAM, sizeof( testProgram ), &testProgram, NULL );
153 test_error( error, "Unable to get kernel program" );
154 if( testProgram != program )
155 {
156 log_error( "ERROR: Kernel program returned invalid value (expected %p, got %p)\n", program, testProgram );
157 return -1;
158 }
159
160 error = clGetKernelInfo( kernel, CL_KERNEL_CONTEXT, sizeof( testContext ), &testContext, NULL );
161 test_error( error, "Unable to get kernel context" );
162 if( testContext != context )
163 {
164 log_error( "ERROR: Kernel context returned invalid value (expected %p, got %p)\n", context, testContext );
165 return -1;
166 }
167
168 /* Release memory */
169 clReleaseKernel( kernel );
170 clReleaseProgram( program );
171 return 0;
172 }
173
test_execute_kernel_local_sizes(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)174 int test_execute_kernel_local_sizes(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
175 {
176 int error;
177 clProgramWrapper program;
178 clKernelWrapper kernel;
179 clMemWrapper streams[2];
180 size_t threads[1], localThreads[1];
181 RandomSeed seed( gRandomSeed );
182 int i;
183
184 num_elements = 100;
185 std::vector<cl_float> inputData(num_elements);
186 std::vector<cl_int> outputData(num_elements);
187
188 /* Create a kernel to test with */
189 if( create_single_kernel_helper( context, &program, &kernel, 1, sample_single_test_kernel, "sample_test" ) != 0 )
190 {
191 return -1;
192 }
193
194 /* Create some I/O streams */
195 streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
196 sizeof(cl_float) * num_elements, NULL, &error);
197 test_error( error, "Creating test array failed" );
198 streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
199 sizeof(cl_int) * num_elements, NULL, &error);
200 test_error( error, "Creating test array failed" );
201
202 /* Write some test data */
203 for (i = 0; i < num_elements; i++)
204 inputData[i] = get_random_float(-(float) 0x7fffffff, (float) 0x7fffffff, seed);
205
206 error = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0,
207 sizeof(cl_float) * num_elements,
208 (void *)inputData.data(), 0, NULL, NULL);
209 test_error( error, "Unable to set testing kernel data" );
210
211 /* Set the arguments */
212 error = clSetKernelArg( kernel, 0, sizeof( streams[0] ), &streams[0] );
213 test_error( error, "Unable to set kernel arguments" );
214 error = clSetKernelArg( kernel, 1, sizeof( streams[1] ), &streams[1] );
215 test_error( error, "Unable to set kernel arguments" );
216
217 /* Test running the kernel and verifying it */
218 threads[0] = (size_t)num_elements;
219 error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
220 test_error( error, "Unable to get work group size to use" );
221
222 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
223 test_error( error, "Kernel execution failed" );
224
225 error = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0,
226 sizeof(cl_int) * num_elements,
227 (void *)outputData.data(), 0, NULL, NULL);
228 test_error( error, "Unable to get result data" );
229
230 for (i = 0; i < num_elements; i++)
231 {
232 if (outputData[i] != (int)inputData[i])
233 {
234 log_error( "ERROR: Data did not verify on first pass!\n" );
235 return -1;
236 }
237 }
238
239 /* Try again */
240 if( localThreads[0] > 1 )
241 localThreads[0] /= 2;
242 while( localThreads[0] > 1 && 0 != threads[0] % localThreads[0] )
243 localThreads[0]--;
244 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
245 test_error( error, "Kernel execution failed" );
246
247 error = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0,
248 sizeof(cl_int) * num_elements,
249 (void *)outputData.data(), 0, NULL, NULL);
250 test_error( error, "Unable to get result data" );
251
252 for (i = 0; i < num_elements; i++)
253 {
254 if (outputData[i] != (int)inputData[i])
255 {
256 log_error( "ERROR: Data did not verify on first pass!\n" );
257 return -1;
258 }
259 }
260
261 /* And again */
262 if( localThreads[0] > 1 )
263 localThreads[0] /= 2;
264 while( localThreads[0] > 1 && 0 != threads[0] % localThreads[0] )
265 localThreads[0]--;
266 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
267 test_error( error, "Kernel execution failed" );
268
269 error = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0,
270 sizeof(cl_int) * num_elements,
271 (void *)outputData.data(), 0, NULL, NULL);
272 test_error( error, "Unable to get result data" );
273
274 for (i = 0; i < num_elements; i++)
275 {
276 if (outputData[i] != (int)inputData[i])
277 {
278 log_error( "ERROR: Data did not verify on first pass!\n" );
279 return -1;
280 }
281 }
282
283 /* One more time */
284 localThreads[0] = (unsigned int)1;
285 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
286 test_error( error, "Kernel execution failed" );
287
288 error = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0,
289 sizeof(cl_int) * num_elements,
290 (void *)outputData.data(), 0, NULL, NULL);
291 test_error( error, "Unable to get result data" );
292
293 for (i = 0; i < num_elements; i++)
294 {
295 if (outputData[i] != (int)inputData[i])
296 {
297 log_error( "ERROR: Data did not verify on first pass!\n" );
298 return -1;
299 }
300 }
301
302 return 0;
303 }
304
test_set_kernel_arg_by_index(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)305 int test_set_kernel_arg_by_index(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
306 {
307 int error;
308 clProgramWrapper program;
309 clKernelWrapper kernel;
310 clMemWrapper streams[2];
311 size_t threads[1], localThreads[1];
312 RandomSeed seed( gRandomSeed );
313 int i;
314
315 num_elements = 10;
316 std::vector<cl_float> inputData(num_elements);
317 std::vector<cl_int> outputData(num_elements);
318
319 /* Create a kernel to test with */
320 if( create_single_kernel_helper( context, &program, &kernel, 1, sample_single_test_kernel, "sample_test" ) != 0 )
321 {
322 return -1;
323 }
324
325 /* Create some I/O streams */
326 streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
327 sizeof(cl_float) * num_elements, NULL, &error);
328 test_error( error, "Creating test array failed" );
329 streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
330 sizeof(cl_int) * num_elements, NULL, &error);
331 test_error( error, "Creating test array failed" );
332
333 /* Write some test data */
334 for (i = 0; i < num_elements; i++)
335 inputData[i] = get_random_float(-(float) 0x7fffffff, (float) 0x7fffffff, seed);
336
337 error = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0,
338 sizeof(cl_float) * num_elements,
339 (void *)inputData.data(), 0, NULL, NULL);
340 test_error( error, "Unable to set testing kernel data" );
341
342 /* Test setting the arguments by index manually */
343 error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1]);
344 test_error( error, "Unable to set indexed kernel arguments" );
345 error = clSetKernelArg(kernel, 0, sizeof( streams[0] ), &streams[0]);
346 test_error( error, "Unable to set indexed kernel arguments" );
347
348
349 /* Test running the kernel and verifying it */
350 threads[0] = (size_t)num_elements;
351
352 error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
353 test_error( error, "Unable to get work group size to use" );
354
355 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
356 test_error( error, "Kernel execution failed" );
357
358 error = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0,
359 sizeof(cl_int) * num_elements,
360 (void *)outputData.data(), 0, NULL, NULL);
361 test_error( error, "Unable to get result data" );
362
363 for (i = 0; i < num_elements; i++)
364 {
365 if (outputData[i] != (int)inputData[i])
366 {
367 log_error( "ERROR: Data did not verify on first pass!\n" );
368 return -1;
369 }
370 }
371
372 return 0;
373 }
374
test_set_kernel_arg_constant(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)375 int test_set_kernel_arg_constant(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
376 {
377 int error;
378 clProgramWrapper program;
379 clKernelWrapper kernel;
380 clMemWrapper streams[3];
381 size_t threads[1], localThreads[1];
382 int i;
383 cl_ulong maxSize;
384 MTdata d;
385
386 num_elements = 10;
387 std::vector<cl_int> outputData(num_elements);
388 std::vector<cl_int> randomTestDataA(num_elements);
389 std::vector<cl_int> randomTestDataB(num_elements);
390
391 /* Verify our test buffer won't be bigger than allowed */
392 error = clGetDeviceInfo( deviceID, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof( maxSize ), &maxSize, 0 );
393 test_error( error, "Unable to get max constant buffer size" );
394 if (maxSize < sizeof(cl_int) * num_elements)
395 {
396 log_error( "ERROR: Unable to test constant argument to kernel: max size of constant buffer is reported as %d!\n", (int)maxSize );
397 return -1;
398 }
399
400 /* Create a kernel to test with */
401 if( create_single_kernel_helper( context, &program, &kernel, 1, sample_const_test_kernel, "sample_test" ) != 0 )
402 {
403 return -1;
404 }
405
406 /* Create some I/O streams */
407 d = init_genrand( gRandomSeed );
408 for (i = 0; i < num_elements; i++)
409 {
410 randomTestDataA[i] = (cl_int)genrand_int32(d) & 0xffffff; /* Make sure values are positive, just so we don't have to */
411 randomTestDataB[i] = (cl_int)genrand_int32(d) & 0xffffff; /* deal with overflow on the verification */
412 }
413 free_mtdata(d); d = NULL;
414
415 streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
416 sizeof(cl_int) * num_elements,
417 randomTestDataA.data(), &error);
418 test_error( error, "Creating test array failed" );
419 streams[1] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
420 sizeof(cl_int) * num_elements,
421 randomTestDataB.data(), &error);
422 test_error( error, "Creating test array failed" );
423 streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
424 sizeof(cl_int) * num_elements, NULL, &error);
425 test_error( error, "Creating test array failed" );
426
427 /* Set the arguments */
428 error = clSetKernelArg(kernel, 0, sizeof( streams[0] ), &streams[0]);
429 test_error( error, "Unable to set indexed kernel arguments" );
430 error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1]);
431 test_error( error, "Unable to set indexed kernel arguments" );
432 error = clSetKernelArg(kernel, 2, sizeof( streams[2] ), &streams[2]);
433 test_error( error, "Unable to set indexed kernel arguments" );
434
435
436 /* Test running the kernel and verifying it */
437 threads[0] = (size_t)num_elements;
438
439 error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
440 test_error( error, "Unable to get work group size to use" );
441
442 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
443 test_error( error, "Kernel execution failed" );
444
445 error = clEnqueueReadBuffer(queue, streams[2], CL_TRUE, 0,
446 sizeof(cl_int) * num_elements,
447 (void *)outputData.data(), 0, NULL, NULL);
448 test_error( error, "Unable to get result data" );
449
450 for (i = 0; i < num_elements; i++)
451 {
452 if (outputData[i] != randomTestDataA[i] + randomTestDataB[i])
453 {
454 log_error( "ERROR: Data sample %d did not verify! %d does not match %d + %d (%d)\n", i, outputData[i], randomTestDataA[i], randomTestDataB[i], ( randomTestDataA[i] + randomTestDataB[i] ) );
455 return -1;
456 }
457 }
458
459 return 0;
460 }
461
test_set_kernel_arg_struct_array(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)462 int test_set_kernel_arg_struct_array(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
463 {
464 int error;
465 clProgramWrapper program;
466 clKernelWrapper kernel;
467 clMemWrapper streams[2];
468 size_t threads[1], localThreads[1];
469 int i;
470 MTdata d;
471
472 num_elements = 10;
473 std::vector<cl_int> outputData(num_elements);
474
475 typedef struct img_pair_type
476 {
477 int A;
478 int B;
479 } image_pair_t;
480
481 std::vector<image_pair_t> image_pair(num_elements);
482
483
484 /* Create a kernel to test with */
485 if( create_single_kernel_helper( context, &program, &kernel, 1, sample_struct_array_test_kernel, "sample_test" ) != 0 )
486 {
487 return -1;
488 }
489
490 /* Create some I/O streams */
491 d = init_genrand( gRandomSeed );
492 for (i = 0; i < num_elements; i++)
493 {
494 image_pair[i].A = (cl_int)genrand_int32(d);
495 image_pair[i].B = (cl_int)genrand_int32(d);
496 }
497 free_mtdata(d); d = NULL;
498
499 streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
500 sizeof(image_pair_t) * num_elements,
501 (void *)image_pair.data(), &error);
502 test_error( error, "Creating test array failed" );
503 streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
504 sizeof(cl_int) * num_elements, NULL, &error);
505 test_error( error, "Creating test array failed" );
506
507 /* Set the arguments */
508 error = clSetKernelArg(kernel, 0, sizeof( streams[0] ), &streams[0]);
509 test_error( error, "Unable to set indexed kernel arguments" );
510 error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1]);
511 test_error( error, "Unable to set indexed kernel arguments" );
512
513 /* Test running the kernel and verifying it */
514 threads[0] = (size_t)num_elements;
515
516 error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
517 test_error( error, "Unable to get work group size to use" );
518
519 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
520 test_error( error, "Kernel execution failed" );
521
522 error = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0,
523 sizeof(cl_int) * num_elements,
524 (void *)outputData.data(), 0, NULL, NULL);
525 test_error( error, "Unable to get result data" );
526
527 for (i = 0; i < num_elements; i++)
528 {
529 if (outputData[i] != image_pair[i].A + image_pair[i].B)
530 {
531 log_error( "ERROR: Data did not verify!\n" );
532 return -1;
533 }
534 }
535
536 return 0;
537 }
538
test_create_kernels_in_program(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)539 int test_create_kernels_in_program(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
540 {
541 int error;
542 cl_program program;
543 cl_kernel kernel[3];
544 unsigned int kernelCount;
545
546 error = create_single_kernel_helper(context, &program, NULL, 2, sample_two_kernel_program, NULL);
547 test_error(error, "Unable to build test program");
548
549 /* Try getting the kernel count */
550 error = clCreateKernelsInProgram( program, 0, NULL, &kernelCount );
551 test_error( error, "Unable to get kernel count for built program" );
552 if( kernelCount != 2 )
553 {
554 log_error( "ERROR: Returned kernel count from clCreateKernelsInProgram is incorrect! (got %d, expected 2)\n", kernelCount );
555 return -1;
556 }
557
558 /* Try actually getting the kernels */
559 error = clCreateKernelsInProgram( program, 2, kernel, NULL );
560 test_error( error, "Unable to get kernels for built program" );
561 clReleaseKernel( kernel[0] );
562 clReleaseKernel( kernel[1] );
563
564 clReleaseProgram( program );
565 return 0;
566 }
567
test_kernel_global_constant(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)568 int test_kernel_global_constant(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
569 {
570 int error;
571 clProgramWrapper program;
572 clKernelWrapper kernel;
573 clMemWrapper streams[2];
574 size_t threads[1], localThreads[1];
575 int i;
576 MTdata d;
577
578 num_elements = 10;
579 std::vector<cl_int> outputData(num_elements);
580 std::vector<cl_int> randomTestDataA(num_elements);
581
582 /* Create a kernel to test with */
583 if( create_single_kernel_helper( context, &program, &kernel, 1, sample_const_global_test_kernel, "sample_test" ) != 0 )
584 {
585 return -1;
586 }
587
588 /* Create some I/O streams */
589 d = init_genrand( gRandomSeed );
590 for (i = 0; i < num_elements; i++)
591 {
592 randomTestDataA[i] = (cl_int)genrand_int32(d) & 0xffff; /* Make sure values are positive and small, just so we don't have to */
593 }
594 free_mtdata(d); d = NULL;
595
596 streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
597 sizeof(cl_int) * num_elements,
598 randomTestDataA.data(), &error);
599 test_error( error, "Creating test array failed" );
600 streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
601 sizeof(cl_int) * num_elements, NULL, &error);
602 test_error( error, "Creating test array failed" );
603
604 /* Set the arguments */
605 error = clSetKernelArg(kernel, 0, sizeof( streams[0] ), &streams[0]);
606 test_error( error, "Unable to set indexed kernel arguments" );
607 error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1]);
608 test_error( error, "Unable to set indexed kernel arguments" );
609
610
611 /* Test running the kernel and verifying it */
612 threads[0] = (size_t)num_elements;
613
614 error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
615 test_error( error, "Unable to get work group size to use" );
616
617 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
618 test_error( error, "Kernel execution failed" );
619
620 error = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0,
621 sizeof(cl_int) * num_elements,
622 (void *)outputData.data(), 0, NULL, NULL);
623 test_error( error, "Unable to get result data" );
624
625 for (i = 0; i < num_elements; i++)
626 {
627 if (outputData[i] != randomTestDataA[i] + 1024)
628 {
629 log_error( "ERROR: Data sample %d did not verify! %d does not match %d + 1024 (%d)\n", i, outputData[i], randomTestDataA[i], ( randomTestDataA[i] + 1024 ) );
630 return -1;
631 }
632 }
633
634 return 0;
635 }
636
637
638
639