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 "harness/compat.h"
17
18 #include <stdio.h>
19 #include <stdlib.h>
20 #include <string.h>
21 #include <sys/types.h>
22 #include <sys/stat.h>
23
24 #include "procs.h"
25 #include "harness/errorHelpers.h"
26
27
28 #ifndef uchar
29 typedef unsigned char uchar;
30 #endif
31
32 #ifndef TestStruct
33 typedef struct{
34 int a;
35 float b;
36 } TestStruct;
37 #endif
38
39 // If this is set to 1 the writes are done via map/unmap
40 static int gTestMap = 0;
41
42 const char *buffer_write_int_kernel_code[] = {
43 "__kernel void test_buffer_write_int(__global int *src, __global int *dst)\n"
44 "{\n"
45 " int tid = get_global_id(0);\n"
46 "\n"
47 " dst[tid] = src[tid];\n"
48 "}\n",
49
50 "__kernel void test_buffer_write_int2(__global int2 *src, __global int2 *dst)\n"
51 "{\n"
52 " int tid = get_global_id(0);\n"
53 "\n"
54 " dst[tid] = src[tid];\n"
55 "}\n",
56
57 "__kernel void test_buffer_write_int4(__global int4 *src, __global int4 *dst)\n"
58 "{\n"
59 " int tid = get_global_id(0);\n"
60 "\n"
61 " dst[tid] = src[tid];\n"
62 "}\n",
63
64 "__kernel void test_buffer_write_int8(__global int8 *src, __global int8 *dst)\n"
65 "{\n"
66 " int tid = get_global_id(0);\n"
67 "\n"
68 " dst[tid] = src[tid];\n"
69 "}\n",
70
71 "__kernel void test_buffer_write_int16(__global int16 *src, __global int16 *dst)\n"
72 "{\n"
73 " int tid = get_global_id(0);\n"
74 "\n"
75 " dst[tid] = src[tid];\n"
76 "}\n" };
77
78 static const char *int_kernel_name[] = { "test_buffer_write_int", "test_buffer_write_int2", "test_buffer_write_int4", "test_buffer_write_int8", "test_buffer_write_int16" };
79
80
81 const char *buffer_write_uint_kernel_code[] = {
82 "__kernel void test_buffer_write_uint(__global uint *src, __global uint *dst)\n"
83 "{\n"
84 " int tid = get_global_id(0);\n"
85 "\n"
86 " dst[tid] = src[tid];\n"
87 "}\n",
88
89 "__kernel void test_buffer_write_uint2(__global uint2 *src, __global uint2 *dst)\n"
90 "{\n"
91 " int tid = get_global_id(0);\n"
92 "\n"
93 " dst[tid] = src[tid];\n"
94 "}\n",
95
96 "__kernel void test_buffer_write_uint4(__global uint4 *src, __global uint4 *dst)\n"
97 "{\n"
98 " int tid = get_global_id(0);\n"
99 "\n"
100 " dst[tid] = src[tid];\n"
101 "}\n",
102
103 "__kernel void test_buffer_write_uint8(__global uint8 *src, __global uint8 *dst)\n"
104 "{\n"
105 " int tid = get_global_id(0);\n"
106 "\n"
107 " dst[tid] = src[tid];\n"
108 "}\n",
109
110 "__kernel void test_buffer_write_uint16(__global uint16 *src, __global uint16 *dst)\n"
111 "{\n"
112 " int tid = get_global_id(0);\n"
113 "\n"
114 " dst[tid] = src[tid];\n"
115 "}\n" };
116
117 static const char *uint_kernel_name[] = { "test_buffer_write_uint", "test_buffer_write_uint2", "test_buffer_write_uint4", "test_buffer_write_uint8", "test_buffer_write_uint16" };
118
119
120 const char *buffer_write_ushort_kernel_code[] = {
121 "__kernel void test_buffer_write_ushort(__global ushort *src, __global ushort *dst)\n"
122 "{\n"
123 " int tid = get_global_id(0);\n"
124 "\n"
125 " dst[tid] = src[tid];\n"
126 "}\n",
127
128 "__kernel void test_buffer_write_ushort2(__global ushort2 *src, __global ushort2 *dst)\n"
129 "{\n"
130 " int tid = get_global_id(0);\n"
131 "\n"
132 " dst[tid] = src[tid];\n"
133 "}\n",
134
135 "__kernel void test_buffer_write_ushort4(__global ushort4 *src, __global ushort4 *dst)\n"
136 "{\n"
137 " int tid = get_global_id(0);\n"
138 "\n"
139 " dst[tid] = src[tid];\n"
140 "}\n",
141
142 "__kernel void test_buffer_write_ushort8(__global ushort8 *src, __global ushort8 *dst)\n"
143 "{\n"
144 " int tid = get_global_id(0);\n"
145 "\n"
146 " dst[tid] = src[tid];\n"
147 "}\n",
148
149 "__kernel void test_buffer_write_ushort16(__global ushort16 *src, __global ushort16 *dst)\n"
150 "{\n"
151 " int tid = get_global_id(0);\n"
152 "\n"
153 " dst[tid] = src[tid];\n"
154 "}\n" };
155
156 static const char *ushort_kernel_name[] = { "test_buffer_write_ushort", "test_buffer_write_ushort2", "test_buffer_write_ushort4", "test_buffer_write_ushort8", "test_buffer_write_ushort16" };
157
158
159
160 const char *buffer_write_short_kernel_code[] = {
161 "__kernel void test_buffer_write_short(__global short *src, __global short *dst)\n"
162 "{\n"
163 " int tid = get_global_id(0);\n"
164 "\n"
165 " dst[tid] = src[tid];\n"
166 "}\n",
167
168 "__kernel void test_buffer_write_short2(__global short2 *src, __global short2 *dst)\n"
169 "{\n"
170 " int tid = get_global_id(0);\n"
171 "\n"
172 " dst[tid] = src[tid];\n"
173 "}\n",
174
175 "__kernel void test_buffer_write_short4(__global short4 *src, __global short4 *dst)\n"
176 "{\n"
177 " int tid = get_global_id(0);\n"
178 "\n"
179 " dst[tid] = src[tid];\n"
180 "}\n",
181
182 "__kernel void test_buffer_write_short8(__global short8 *src, __global short8 *dst)\n"
183 "{\n"
184 " int tid = get_global_id(0);\n"
185 "\n"
186 " dst[tid] = src[tid];\n"
187 "}\n",
188
189 "__kernel void test_buffer_write_short16(__global short16 *src, __global short16 *dst)\n"
190 "{\n"
191 " int tid = get_global_id(0);\n"
192 "\n"
193 " dst[tid] = src[tid];\n"
194 "}\n" };
195
196 static const char *short_kernel_name[] = { "test_buffer_write_short", "test_buffer_write_short2", "test_buffer_write_short4", "test_buffer_write_short8", "test_buffer_write_short16" };
197
198
199 const char *buffer_write_char_kernel_code[] = {
200 "__kernel void test_buffer_write_char(__global char *src, __global char *dst)\n"
201 "{\n"
202 " int tid = get_global_id(0);\n"
203 "\n"
204 " dst[tid] = src[tid];\n"
205 "}\n",
206
207 "__kernel void test_buffer_write_char2(__global char2 *src, __global char2 *dst)\n"
208 "{\n"
209 " int tid = get_global_id(0);\n"
210 "\n"
211 " dst[tid] = src[tid];\n"
212 "}\n",
213
214 "__kernel void test_buffer_write_char4(__global char4 *src, __global char4 *dst)\n"
215 "{\n"
216 " int tid = get_global_id(0);\n"
217 "\n"
218 " dst[tid] = src[tid];\n"
219 "}\n",
220
221 "__kernel void test_buffer_write_char8(__global char8 *src, __global char8 *dst)\n"
222 "{\n"
223 " int tid = get_global_id(0);\n"
224 "\n"
225 " dst[tid] = src[tid];\n"
226 "}\n",
227
228 "__kernel void test_buffer_write_char16(__global char16 *src, __global char16 *dst)\n"
229 "{\n"
230 " int tid = get_global_id(0);\n"
231 "\n"
232 " dst[tid] = src[tid];\n"
233 "}\n" };
234
235 static const char *char_kernel_name[] = { "test_buffer_write_char", "test_buffer_write_char2", "test_buffer_write_char4", "test_buffer_write_char8", "test_buffer_write_char16" };
236
237
238 const char *buffer_write_uchar_kernel_code[] = {
239 "__kernel void test_buffer_write_uchar(__global uchar *src, __global uchar *dst)\n"
240 "{\n"
241 " int tid = get_global_id(0);\n"
242 "\n"
243 " dst[tid] = src[tid];\n"
244 "}\n",
245
246 "__kernel void test_buffer_write_uchar2(__global uchar2 *src, __global uchar2 *dst)\n"
247 "{\n"
248 " int tid = get_global_id(0);\n"
249 "\n"
250 " dst[tid] = src[tid];\n"
251 "}\n",
252
253 "__kernel void test_buffer_write_uchar4(__global uchar4 *src, __global uchar4 *dst)\n"
254 "{\n"
255 " int tid = get_global_id(0);\n"
256 "\n"
257 " dst[tid] = src[tid];\n"
258 "}\n",
259
260 "__kernel void test_buffer_write_uchar8(__global uchar8 *src, __global uchar8 *dst)\n"
261 "{\n"
262 " int tid = get_global_id(0);\n"
263 "\n"
264 " dst[tid] = src[tid];\n"
265 "}\n",
266
267 "__kernel void test_buffer_write_uchar16(__global uchar16 *src, __global uchar16 *dst)\n"
268 "{\n"
269 " int tid = get_global_id(0);\n"
270 "\n"
271 " dst[tid] = src[tid];\n"
272 "}\n" };
273
274 static const char *uchar_kernel_name[] = { "test_buffer_write_uchar", "test_buffer_write_uchar2", "test_buffer_write_uchar4", "test_buffer_write_uchar8", "test_buffer_write_uchar16" };
275
276
277 const char *buffer_write_float_kernel_code[] = {
278 "__kernel void test_buffer_write_float(__global float *src, __global float *dst)\n"
279 "{\n"
280 " int tid = get_global_id(0);\n"
281 "\n"
282 " dst[tid] = src[tid];\n"
283 "}\n",
284
285 "__kernel void test_buffer_write_float2(__global float2 *src, __global float2 *dst)\n"
286 "{\n"
287 " int tid = get_global_id(0);\n"
288 "\n"
289 " dst[tid] = src[tid];\n"
290 "}\n",
291
292 "__kernel void test_buffer_write_float4(__global float4 *src, __global float4 *dst)\n"
293 "{\n"
294 " int tid = get_global_id(0);\n"
295 "\n"
296 " dst[tid] = src[tid];\n"
297 "}\n",
298
299 "__kernel void test_buffer_write_float8(__global float8 *src, __global float8 *dst)\n"
300 "{\n"
301 " int tid = get_global_id(0);\n"
302 "\n"
303 " dst[tid] = src[tid];\n"
304 "}\n",
305
306 "__kernel void test_buffer_write_float16(__global float16 *src, __global float16 *dst)\n"
307 "{\n"
308 " int tid = get_global_id(0);\n"
309 "\n"
310 " dst[tid] = src[tid];\n"
311 "}\n" };
312
313 static const char *float_kernel_name[] = { "test_buffer_write_float", "test_buffer_write_float2", "test_buffer_write_float4", "test_buffer_write_float8", "test_buffer_write_float16" };
314
315
316 const char *buffer_write_half_kernel_code[] = {
317 "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
318 "__kernel void test_buffer_write_half(__global half *src, __global half "
319 "*dst)\n"
320 "{\n"
321 " int tid = get_global_id(0);\n"
322 "\n"
323 " dst[tid] = src[tid];\n"
324 "}\n",
325
326 "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
327 "__kernel void test_buffer_write_half2(__global half2 *src, __global half2 "
328 "*dst)\n"
329 "{\n"
330 " int tid = get_global_id(0);\n"
331 "\n"
332 " dst[tid] = src[tid];\n"
333 "}\n",
334
335 "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
336 "__kernel void test_buffer_write_half4(__global half4 *src, __global half4 "
337 "*dst)\n"
338 "{\n"
339 " int tid = get_global_id(0);\n"
340 "\n"
341 " dst[tid] = src[tid];\n"
342 "}\n",
343
344 "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
345 "__kernel void test_buffer_write_half8(__global half8 *src, __global half8 "
346 "*dst)\n"
347 "{\n"
348 " int tid = get_global_id(0);\n"
349 "\n"
350 " dst[tid] = src[tid];\n"
351 "}\n",
352
353 "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
354 "__kernel void test_buffer_write_half16(__global half16 *src, __global "
355 "half16 *dst)\n"
356 "{\n"
357 " int tid = get_global_id(0);\n"
358 "\n"
359 " dst[tid] = src[tid];\n"
360 "}\n"
361 };
362
363 static const char *half_kernel_name[] = { "test_buffer_write_half", "test_buffer_write_half2", "test_buffer_write_half4", "test_buffer_write_half8", "test_buffer_write_half16" };
364
365
366 const char *buffer_write_long_kernel_code[] = {
367 "__kernel void test_buffer_write_long(__global long *src, __global long *dst)\n"
368 "{\n"
369 " int tid = get_global_id(0);\n"
370 "\n"
371 " dst[tid] = src[tid];\n"
372 "}\n",
373
374 "__kernel void test_buffer_write_long2(__global long2 *src, __global long2 *dst)\n"
375 "{\n"
376 " int tid = get_global_id(0);\n"
377 "\n"
378 " dst[tid] = src[tid];\n"
379 "}\n",
380
381 "__kernel void test_buffer_write_long4(__global long4 *src, __global long4 *dst)\n"
382 "{\n"
383 " int tid = get_global_id(0);\n"
384 "\n"
385 " dst[tid] = src[tid];\n"
386 "}\n",
387
388 "__kernel void test_buffer_write_long8(__global long8 *src, __global long8 *dst)\n"
389 "{\n"
390 " int tid = get_global_id(0);\n"
391 "\n"
392 " dst[tid] = src[tid];\n"
393 "}\n",
394
395 "__kernel void test_buffer_write_long16(__global long16 *src, __global long16 *dst)\n"
396 "{\n"
397 " int tid = get_global_id(0);\n"
398 "\n"
399 " dst[tid] = src[tid];\n"
400 "}\n" };
401
402 static const char *long_kernel_name[] = { "test_buffer_write_long", "test_buffer_write_long2", "test_buffer_write_long4", "test_buffer_write_long8", "test_buffer_write_long16" };
403
404
405 const char *buffer_write_ulong_kernel_code[] = {
406 "__kernel void test_buffer_write_ulong(__global ulong *src, __global ulong *dst)\n"
407 "{\n"
408 " int tid = get_global_id(0);\n"
409 "\n"
410 " dst[tid] = src[tid];\n"
411 "}\n",
412
413 "__kernel void test_buffer_write_ulong2(__global ulong2 *src, __global ulong2 *dst)\n"
414 "{\n"
415 " int tid = get_global_id(0);\n"
416 "\n"
417 " dst[tid] = src[tid];\n"
418 "}\n",
419
420 "__kernel void test_buffer_write_ulong4(__global ulong4 *src, __global ulong4 *dst)\n"
421 "{\n"
422 " int tid = get_global_id(0);\n"
423 "\n"
424 " dst[tid] = src[tid];\n"
425 "}\n",
426
427 "__kernel void test_buffer_write_ulong8(__global ulong8 *src, __global ulong8 *dst)\n"
428 "{\n"
429 " int tid = get_global_id(0);\n"
430 "\n"
431 " dst[tid] = src[tid];\n"
432 "}\n",
433
434 "__kernel void test_buffer_write_ulong16(__global ulong16 *src, __global ulong16 *dst)\n"
435 "{\n"
436 " int tid = get_global_id(0);\n"
437 "\n"
438 " dst[tid] = src[tid];\n"
439 "}\n" };
440
441 static const char *ulong_kernel_name[] = { "test_buffer_write_ulong", "test_buffer_write_ulong2", "test_buffer_write_ulong4", "test_buffer_write_ulong8", "test_buffer_write_ulong16" };
442
443
444 static const char *struct_kernel_code =
445 "typedef struct{\n"
446 "int a;\n"
447 "float b;\n"
448 "} TestStruct;\n"
449 "__kernel void read_write_struct(__global TestStruct *src, __global TestStruct *dst)\n"
450 "{\n"
451 " int tid = get_global_id(0);\n"
452 "\n"
453 " dst[tid].a = src[tid].a;\n"
454 " dst[tid].b = src[tid].b;\n"
455 "}\n";
456
457
458
verify_write_int(void * ptr1,void * ptr2,int n)459 static int verify_write_int( void *ptr1, void *ptr2, int n )
460 {
461 int i;
462 int *inptr = (int *)ptr1;
463 int *outptr = (int *)ptr2;
464
465 for (i=0; i<n; i++){
466 if ( outptr[i] != inptr[i] )
467 return -1;
468 }
469
470 return 0;
471 }
472
473
verify_write_uint(void * ptr1,void * ptr2,int n)474 static int verify_write_uint( void *ptr1, void *ptr2, int n )
475 {
476 int i;
477 cl_uint *inptr = (cl_uint *)ptr1;
478 cl_uint *outptr = (cl_uint *)ptr2;
479
480 for (i=0; i<n; i++){
481 if ( outptr[i] != inptr[i] )
482 return -1;
483 }
484
485 return 0;
486 }
487
488
verify_write_short(void * ptr1,void * ptr2,int n)489 static int verify_write_short( void *ptr1, void *ptr2, int n )
490 {
491 int i;
492 short *inptr = (short *)ptr1;
493 short *outptr = (short *)ptr2;
494
495 for (i=0; i<n; i++){
496 if ( outptr[i] != inptr[i] )
497 return -1;
498 }
499
500 return 0;
501 }
502
503
verify_write_ushort(void * ptr1,void * ptr2,int n)504 static int verify_write_ushort( void *ptr1, void *ptr2, int n )
505 {
506 int i;
507 cl_ushort *inptr = (cl_ushort *)ptr1;
508 cl_ushort *outptr = (cl_ushort *)ptr2;
509
510 for (i=0; i<n; i++){
511 if ( outptr[i] != inptr[i] )
512 return -1;
513 }
514
515 return 0;
516 }
517
518
verify_write_char(void * ptr1,void * ptr2,int n)519 static int verify_write_char( void *ptr1, void *ptr2, int n )
520 {
521 int i;
522 char *inptr = (char *)ptr1;
523 char *outptr = (char *)ptr2;
524
525 for (i=0; i<n; i++){
526 if ( outptr[i] != inptr[i] )
527 return -1;
528 }
529
530 return 0;
531 }
532
533
verify_write_uchar(void * ptr1,void * ptr2,int n)534 static int verify_write_uchar( void *ptr1, void *ptr2, int n )
535 {
536 int i;
537 uchar *inptr = (uchar *)ptr1;
538 uchar *outptr = (uchar *)ptr2;
539
540 for (i=0; i<n; i++){
541 if ( outptr[i] != inptr[i] )
542 return -1;
543 }
544
545 return 0;
546 }
547
548
verify_write_float(void * ptr1,void * ptr2,int n)549 static int verify_write_float( void *ptr1, void *ptr2, int n )
550 {
551 int i;
552 float *inptr = (float *)ptr1;
553 float *outptr = (float *)ptr2;
554
555 for (i=0; i<n; i++){
556 if ( outptr[i] != inptr[i] )
557 return -1;
558 }
559
560 return 0;
561 }
562
563
verify_write_half(void * ptr1,void * ptr2,int n)564 static int verify_write_half( void *ptr1, void *ptr2, int n )
565 {
566 int i;
567 cl_half *inptr = (cl_half *)ptr1;
568 cl_half *outptr = (cl_half *)ptr2;
569
570 for ( i = 0; i < n; i++ ){
571 if ( outptr[i] != inptr[i] )
572 return -1;
573 }
574
575 return 0;
576 }
577
578
verify_write_long(void * ptr1,void * ptr2,int n)579 static int verify_write_long( void *ptr1, void *ptr2, int n )
580 {
581 int i;
582 cl_long *inptr = (cl_long *)ptr1;
583 cl_long *outptr = (cl_long *)ptr2;
584
585 for (i=0; i<n; i++){
586 if ( outptr[i] != inptr[i] )
587 return -1;
588 }
589
590 return 0;
591 }
592
593
verify_write_ulong(void * ptr1,void * ptr2,int n)594 static int verify_write_ulong( void *ptr1, void *ptr2, int n )
595 {
596 int i;
597 cl_ulong *inptr = (cl_ulong *)ptr1;
598 cl_ulong *outptr = (cl_ulong *)ptr2;
599
600 for (i=0; i<n; i++){
601 if ( outptr[i] != inptr[i] )
602 return -1;
603 }
604
605 return 0;
606 }
607
608
verify_write_struct(void * ptr1,void * ptr2,int n)609 static int verify_write_struct( void *ptr1, void *ptr2, int n )
610 {
611 int i;
612 TestStruct *inptr = (TestStruct *)ptr1;
613 TestStruct *outptr = (TestStruct *)ptr2;
614
615 for (i=0; i<n; i++){
616 if ( ( outptr[i].a != inptr[i].a ) || ( outptr[i].b != outptr[i].b ) )
617 return -1;
618 }
619
620 return 0;
621 }
622
623
test_buffer_write(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements,size_t size,char * type,int loops,void * inptr[5],const char * kernelCode[],const char * kernelName[],int (* fn)(void *,void *,int),MTdata d)624 int test_buffer_write( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, size_t size, char *type, int loops,
625 void *inptr[5], const char *kernelCode[], const char *kernelName[], int (*fn)(void *,void *,int), MTdata d )
626 {
627 void *outptr[5];
628 clProgramWrapper program[5];
629 clKernelWrapper kernel[5];
630 size_t ptrSizes[5];
631 size_t global_work_size[3];
632 cl_int err;
633 int i;
634 int src_flag_id, dst_flag_id;
635 int total_errors = 0;
636
637 size_t min_alignment = get_min_alignment(context);
638
639 global_work_size[0] = (size_t)num_elements;
640
641 ptrSizes[0] = size;
642 ptrSizes[1] = ptrSizes[0] << 1;
643 ptrSizes[2] = ptrSizes[1] << 1;
644 ptrSizes[3] = ptrSizes[2] << 1;
645 ptrSizes[4] = ptrSizes[3] << 1;
646
647 loops = (loops < 5 ? loops : 5);
648 for (i = 0; i < loops; i++)
649 {
650 err = create_single_kernel_helper(context, &program[i], &kernel[i], 1,
651 &kernelCode[i], kernelName[i]);
652 if (err)
653 {
654 log_error(" Error creating program for %s\n", type);
655 return -1;
656 }
657
658 for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++)
659 {
660 for (dst_flag_id = 0; dst_flag_id < NUM_FLAGS; dst_flag_id++)
661 {
662 clMemWrapper buffers[2];
663
664 if ((flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) || (flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR))
665 buffers[0] = clCreateBuffer(context, flag_set[src_flag_id],
666 ptrSizes[i] * num_elements,
667 inptr[i], &err);
668 else
669 buffers[0] =
670 clCreateBuffer(context, flag_set[src_flag_id],
671 ptrSizes[i] * num_elements, NULL, &err);
672
673 if (!buffers[0] || err)
674 {
675 align_free( outptr[i] );
676 print_error(err, " clCreateBuffer failed\n" );
677 return -1;
678 }
679 if ( ! strcmp( type, "half" ) ){
680 outptr[i] = align_malloc( ptrSizes[i] * (num_elements * 2 ), min_alignment);
681 if ((flag_set[dst_flag_id] & CL_MEM_USE_HOST_PTR) || (flag_set[dst_flag_id] & CL_MEM_COPY_HOST_PTR))
682 buffers[1] = clCreateBuffer(
683 context, flag_set[dst_flag_id],
684 ptrSizes[i] * 2 * num_elements, outptr[i], &err);
685 else
686 buffers[1] = clCreateBuffer(
687 context, flag_set[dst_flag_id],
688 ptrSizes[i] * 2 * num_elements, NULL, &err);
689 }
690 else{
691 outptr[i] = align_malloc( ptrSizes[i] * num_elements, min_alignment);
692 if ((flag_set[dst_flag_id] & CL_MEM_USE_HOST_PTR) || (flag_set[dst_flag_id] & CL_MEM_COPY_HOST_PTR))
693 buffers[1] = clCreateBuffer(
694 context, flag_set[dst_flag_id],
695 ptrSizes[i] * num_elements, outptr[i], &err);
696 else
697 buffers[1] = clCreateBuffer(
698 context, flag_set[dst_flag_id],
699 ptrSizes[i] * num_elements, NULL, &err);
700 }
701 if ( err ){
702 align_free( outptr[i] );
703 print_error(err, " clCreateBuffer failed\n" );
704 return -1;
705 }
706
707 if (gTestMap) {
708 void *dataPtr;
709 dataPtr = clEnqueueMapBuffer(
710 queue, buffers[0], CL_TRUE, CL_MAP_WRITE, 0,
711 ptrSizes[i] * num_elements, 0, NULL, NULL, &err);
712 if (err) {
713 print_error(err, "clEnqueueMapBuffer failed");
714 align_free( outptr[i] );
715 return -1;
716 }
717
718 memcpy(dataPtr, inptr[i], ptrSizes[i]*num_elements);
719
720 err = clEnqueueUnmapMemObject(queue, buffers[0], dataPtr, 0,
721 NULL, NULL);
722 if (err) {
723 print_error(err, "clEnqueueUnmapMemObject failed");
724 align_free( outptr[i] );
725 return -1;
726 }
727 }
728 else if (!(flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) && !(flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR)) {
729 err = clEnqueueWriteBuffer(queue, buffers[0], CL_TRUE, 0,
730 ptrSizes[i] * num_elements,
731 inptr[i], 0, NULL, NULL);
732 if ( err != CL_SUCCESS ){
733 align_free( outptr[i] );
734 print_error( err, " clWriteBuffer failed" );
735 return -1;
736 }
737 }
738
739 err = clSetKernelArg(kernel[i], 0, sizeof(cl_mem),
740 (void *)&buffers[0]);
741 err |= clSetKernelArg(kernel[i], 1, sizeof(cl_mem),
742 (void *)&buffers[1]);
743 if ( err != CL_SUCCESS ){
744 align_free( outptr[i] );
745 print_error( err, " clSetKernelArg failed" );
746 return -1;
747 }
748
749 err = clEnqueueNDRangeKernel( queue, kernel[i], 1, NULL, global_work_size, NULL, 0, NULL, NULL );
750 if ( err != CL_SUCCESS ){
751 print_error( err, " clEnqueueNDRangeKernel failed" );
752 align_free( outptr[i] );
753 return -1;
754 }
755
756 err = clEnqueueReadBuffer(queue, buffers[1], true, 0,
757 ptrSizes[i] * num_elements, outptr[i],
758 0, NULL, NULL);
759
760 if ( err != CL_SUCCESS ){
761 align_free( outptr[i] );
762 print_error( err, " clEnqueueReadBuffer failed" );
763 return -1;
764 }
765
766 if ( fn( inptr[i], outptr[i], (int)(ptrSizes[i] * (size_t)num_elements / ptrSizes[0]) ) ){
767 log_error(
768 " %s%d test failed. cl_mem_flags src: %s dst: %s\n",
769 type, 1 << i, flag_set_names[src_flag_id],
770 flag_set_names[dst_flag_id]);
771 total_errors++;
772 }
773 else{
774 log_info(
775 " %s%d test passed. cl_mem_flags src: %s dst: %s\n",
776 type, 1 << i, flag_set_names[src_flag_id],
777 flag_set_names[dst_flag_id]);
778 }
779 // cleanup
780 align_free( outptr[i] );
781 }
782 } // dst cl_mem_flag
783 } // src cl_mem_flag
784
785 return total_errors;
786
787 } // end test_buffer_write()
788
789
790
791
test_buffer_write_struct(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)792 int test_buffer_write_struct( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
793 {
794
795 void *outptr[5];
796 TestStruct *inptr[5];
797 clProgramWrapper program[5];
798 clKernelWrapper kernel[5];
799 size_t ptrSizes[5];
800 size_t size = sizeof( TestStruct );
801 size_t global_work_size[3];
802 cl_int err;
803 int i;
804 cl_uint j;
805 int loops = 1; // no vector for structs
806 int src_flag_id, dst_flag_id;
807 int total_errors = 0;
808 MTdata d = init_genrand( gRandomSeed );
809
810 size_t min_alignment = get_min_alignment(context);
811
812 global_work_size[0] = (size_t)num_elements;
813
814 ptrSizes[0] = size;
815 ptrSizes[1] = ptrSizes[0] << 1;
816 ptrSizes[2] = ptrSizes[1] << 1;
817 ptrSizes[3] = ptrSizes[2] << 1;
818 ptrSizes[4] = ptrSizes[3] << 1;
819
820 loops = (loops < 5 ? loops : 5);
821 for (i = 0; i < loops; i++)
822 {
823
824 err = create_single_kernel_helper(context, &program[i], &kernel[i], 1,
825 &struct_kernel_code,
826 "read_write_struct");
827 if (err)
828 {
829 log_error(" Error creating program for struct\n");
830 free_mtdata(d);
831 return -1;
832 }
833
834 for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++)
835 {
836 for (dst_flag_id = 0; dst_flag_id < NUM_FLAGS; dst_flag_id++)
837 {
838 clMemWrapper buffers[2];
839
840 inptr[i] = (TestStruct *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
841
842 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ ){
843 inptr[i][j].a = (int)genrand_int32(d);
844 inptr[i][j].b = get_random_float( -FLT_MAX, FLT_MAX, d );
845 }
846
847 if ((flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) || (flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR))
848 buffers[0] = clCreateBuffer(context, flag_set[src_flag_id],
849 ptrSizes[i] * num_elements,
850 inptr[i], &err);
851 else
852 buffers[0] =
853 clCreateBuffer(context, flag_set[src_flag_id],
854 ptrSizes[i] * num_elements, NULL, &err);
855 if ( err ){
856 align_free( outptr[i] );
857 print_error(err, " clCreateBuffer failed\n" );
858 free_mtdata(d);
859 return -1;
860 }
861 outptr[i] = align_malloc( ptrSizes[i] * num_elements, min_alignment);
862 if ((flag_set[dst_flag_id] & CL_MEM_USE_HOST_PTR) || (flag_set[dst_flag_id] & CL_MEM_COPY_HOST_PTR))
863 buffers[1] = clCreateBuffer(context, flag_set[dst_flag_id],
864 ptrSizes[i] * num_elements,
865 outptr[i], &err);
866 else
867 buffers[1] =
868 clCreateBuffer(context, flag_set[dst_flag_id],
869 ptrSizes[i] * num_elements, NULL, &err);
870 if (!buffers[1] || err)
871 {
872 align_free( outptr[i] );
873 print_error(err, " clCreateBuffer failed\n" );
874 free_mtdata(d);
875 return -1;
876 }
877
878 if (gTestMap) {
879 void *dataPtr;
880 dataPtr = clEnqueueMapBuffer(
881 queue, buffers[0], CL_TRUE, CL_MAP_WRITE, 0,
882 ptrSizes[i] * num_elements, 0, NULL, NULL, &err);
883 if (err) {
884 print_error(err, "clEnqueueMapBuffer failed");
885 align_free( outptr[i] );
886 free_mtdata(d);
887 return -1;
888 }
889
890 memcpy(dataPtr, inptr[i], ptrSizes[i]*num_elements);
891
892 err = clEnqueueUnmapMemObject(queue, buffers[0], dataPtr, 0,
893 NULL, NULL);
894 if (err) {
895 print_error(err, "clEnqueueUnmapMemObject failed");
896 align_free( outptr[i] );
897 free_mtdata(d);
898 return -1;
899 }
900 }
901 else if (!(flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) && !(flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR)) {
902 err = clEnqueueWriteBuffer(queue, buffers[0], CL_TRUE, 0,
903 ptrSizes[i] * num_elements,
904 inptr[i], 0, NULL, NULL);
905 if ( err != CL_SUCCESS ){
906 align_free( outptr[i] );
907 print_error( err, " clWriteBuffer failed" );
908 free_mtdata(d);
909 return -1;
910 }
911 }
912
913 err = clSetKernelArg(kernel[i], 0, sizeof(cl_mem),
914 (void *)&buffers[0]);
915 err |= clSetKernelArg(kernel[i], 1, sizeof(cl_mem),
916 (void *)&buffers[1]);
917 if ( err != CL_SUCCESS ){
918 align_free( outptr[i] );
919 print_error( err, " clSetKernelArg failed" );
920 free_mtdata(d);
921 return -1;
922 }
923
924 err = clEnqueueNDRangeKernel( queue, kernel[i], 1, NULL, global_work_size, NULL, 0, NULL, NULL );
925 if ( err != CL_SUCCESS ){
926 print_error( err, " clEnqueueNDRangeKernel failed" );
927 align_free( outptr[i] );
928 free_mtdata(d);
929 return -1;
930 }
931
932 err = clEnqueueReadBuffer(queue, buffers[1], true, 0,
933 ptrSizes[i] * num_elements, outptr[i],
934 0, NULL, NULL);
935 if ( err != CL_SUCCESS ){
936 align_free( outptr[i] );
937 print_error( err, " clEnqueueReadBuffer failed" );
938 free_mtdata(d);
939 return -1;
940 }
941
942 if ( verify_write_struct( inptr[i], outptr[i], (int)(ptrSizes[i] * (size_t)num_elements / ptrSizes[0]) ) ){
943 log_error(" buffer_WRITE struct%d test failed. "
944 "cl_mem_flags src: %s dst: %s\n",
945 1 << i, flag_set_names[src_flag_id],
946 flag_set_names[dst_flag_id]);
947 total_errors++;
948 }
949 else{
950 log_info(" buffer_WRITE struct%d test passed. cl_mem_flags "
951 "src: %s dst: %s\n",
952 1 << i, flag_set_names[src_flag_id],
953 flag_set_names[dst_flag_id]);
954 }
955 // cleanup
956 align_free( outptr[i] );
957 align_free( (void *)inptr[i] );
958 }
959 } // dst cl_mem_flag
960 } // src cl_mem_flag
961
962 free_mtdata(d);
963
964 return total_errors;
965
966 } // end test_buffer_struct_write()
967
968
test_buffer_write_array_async(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements,size_t size,char * type,int loops,void * inptr[5],const char * kernelCode[],const char * kernelName[],int (* fn)(void *,void *,int))969 int test_buffer_write_array_async( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, size_t size, char *type, int loops,
970 void *inptr[5], const char *kernelCode[], const char *kernelName[], int (*fn)(void *,void *,int) )
971 {
972 cl_mem buffers[10];
973 void *outptr[5];
974 cl_program program[5];
975 cl_kernel kernel[5];
976 cl_event event[2];
977 size_t ptrSizes[5];
978 size_t global_work_size[3];
979 cl_int err;
980 int i, ii;
981 int src_flag_id, dst_flag_id;
982 int total_errors = 0;
983
984 size_t min_alignment = get_min_alignment(context);
985
986 global_work_size[0] = (size_t)num_elements;
987
988 ptrSizes[0] = size;
989 ptrSizes[1] = ptrSizes[0] << 1;
990 ptrSizes[2] = ptrSizes[1] << 1;
991 ptrSizes[3] = ptrSizes[2] << 1;
992 ptrSizes[4] = ptrSizes[3] << 1;
993
994 for (src_flag_id=0; src_flag_id < NUM_FLAGS; src_flag_id++) {
995 for (dst_flag_id=0; dst_flag_id < NUM_FLAGS; dst_flag_id++) {
996 log_info("Testing with cl_mem_flags src: %s dst: %s\n", flag_set_names[src_flag_id], flag_set_names[dst_flag_id]);
997
998 loops = ( loops < 5 ? loops : 5 );
999 for ( i = 0; i < loops; i++ ){
1000 ii = i << 1;
1001 if ((flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) || (flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR))
1002 buffers[ii] = clCreateBuffer(context, flag_set[src_flag_id], ptrSizes[i] * num_elements, inptr[i], &err);
1003 else
1004 buffers[ii] = clCreateBuffer(context, flag_set[src_flag_id], ptrSizes[i] * num_elements, NULL, &err);
1005 if ( !buffers[ii] || err){
1006 print_error(err, "clCreateBuffer failed\n" );
1007 return -1;
1008 }
1009
1010 outptr[i] = align_malloc( ptrSizes[i] * num_elements, min_alignment);
1011 if ((flag_set[dst_flag_id] & CL_MEM_USE_HOST_PTR) || (flag_set[dst_flag_id] & CL_MEM_COPY_HOST_PTR))
1012 buffers[ii+1] = clCreateBuffer(context, flag_set[dst_flag_id], ptrSizes[i] * num_elements, outptr[i], &err);
1013 else
1014 buffers[ii+1] = clCreateBuffer(context, flag_set[dst_flag_id], ptrSizes[i] * num_elements, NULL, &err);
1015 if ( !buffers[ii+1] || err){
1016 print_error(err, "clCreateBuffer failed\n" );
1017 return -1;
1018 }
1019
1020 err = clEnqueueWriteBuffer(queue, buffers[ii], CL_FALSE, 0, ptrSizes[i]*num_elements, inptr[i], 0, NULL, &(event[0]));
1021 if ( err != CL_SUCCESS ){
1022 print_error( err, "clEnqueueWriteBuffer failed" );
1023 return -1;
1024 }
1025
1026 err = create_single_kernel_helper( context, &program[i], &kernel[i], 1, &kernelCode[i], kernelName[i] );
1027 if ( err ){
1028 log_error( " Error creating program for %s\n", type );
1029 clReleaseMemObject( buffers[ii] );
1030 clReleaseMemObject( buffers[ii+1] );
1031 align_free( outptr[i] );
1032 return -1;
1033 }
1034
1035 err = clSetKernelArg( kernel[i], 0, sizeof( cl_mem ), (void *)&buffers[ii] );
1036 err |= clSetKernelArg( kernel[i], 1, sizeof( cl_mem ), (void *)&buffers[ii+1] );
1037 if ( err != CL_SUCCESS ){
1038 print_error( err, "clSetKernelArg failed" );
1039 clReleaseKernel( kernel[i] );
1040 clReleaseProgram( program[i] );
1041 clReleaseMemObject( buffers[ii] );
1042 clReleaseMemObject( buffers[ii+1] );
1043 align_free( outptr[i] );
1044 return -1;
1045 }
1046
1047 err = clWaitForEvents( 1, &(event[0]) );
1048 if ( err != CL_SUCCESS ){
1049 print_error( err, "clWaitForEvents() failed" );
1050 clReleaseKernel( kernel[i] );
1051 clReleaseProgram( program[i] );
1052 clReleaseMemObject( buffers[ii] );
1053 clReleaseMemObject( buffers[ii+1] );
1054 align_free( outptr[i] );
1055 return -1;
1056 }
1057
1058 err = clEnqueueNDRangeKernel( queue, kernel[i], 1, NULL, global_work_size, NULL, 0, NULL, NULL );
1059
1060 if (err != CL_SUCCESS){
1061 print_error( err, "clEnqueueNDRangeKernel failed" );
1062 return -1;
1063 }
1064
1065 err = clEnqueueReadBuffer( queue, buffers[ii+1], false, 0, ptrSizes[i]*num_elements, outptr[i], 0, NULL, &(event[1]) );
1066 if (err != CL_SUCCESS){
1067 print_error( err, "clEnqueueReadBuffer failed" );
1068 return -1;
1069 }
1070
1071 err = clWaitForEvents( 1, &(event[1]) );
1072 if ( err != CL_SUCCESS ){
1073 print_error( err, "clWaitForEvents() failed" );
1074 }
1075
1076 if ( fn( inptr[i], outptr[i], (int)(ptrSizes[i] * (size_t)num_elements / ptrSizes[0]) ) ){
1077 log_error( " %s%d test failed\n", type, 1<<i );
1078 total_errors++;
1079 }
1080 else{
1081 log_info( " %s%d test passed\n", type, 1<<i );
1082 }
1083
1084 // cleanup
1085 clReleaseEvent( event[0] );
1086 clReleaseEvent( event[1] );
1087 clReleaseMemObject( buffers[ii] );
1088 clReleaseMemObject( buffers[ii+1] );
1089 clReleaseKernel( kernel[i] );
1090 clReleaseProgram( program[i] );
1091 align_free( outptr[i] );
1092 }
1093 } // dst cl_mem_flag
1094 } // src cl_mem_flag
1095
1096 return total_errors;
1097
1098 } // end test_buffer_write_array_async()
1099
1100
test_buffer_write_int(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1101 int test_buffer_write_int( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1102 {
1103 int *inptr[5];
1104 size_t ptrSizes[5];
1105 int i, err;
1106 cl_uint j;
1107 int (*foo)(void *,void *,int);
1108 MTdata d = init_genrand( gRandomSeed );
1109
1110 size_t min_alignment = get_min_alignment(context);
1111
1112 foo = verify_write_int;
1113
1114 ptrSizes[0] = sizeof(cl_int);
1115 ptrSizes[1] = ptrSizes[0] << 1;
1116 ptrSizes[2] = ptrSizes[1] << 1;
1117 ptrSizes[3] = ptrSizes[2] << 1;
1118 ptrSizes[4] = ptrSizes[3] << 1;
1119
1120 for ( i = 0; i < 5; i++ ){
1121 inptr[i] = (int *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1122
1123 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1124 inptr[i][j] = (int)genrand_int32(d);
1125 }
1126
1127 err = test_buffer_write( deviceID, context, queue, num_elements, sizeof( cl_int ), (char*)"int", 5, (void**)inptr,
1128 buffer_write_int_kernel_code, int_kernel_name, foo, d );
1129
1130 for ( i = 0; i < 5; i++ ){
1131 align_free( (void *)inptr[i] );
1132 }
1133 free_mtdata(d);
1134
1135 return err;
1136
1137 } // end test_buffer_int_write()
1138
1139
test_buffer_write_uint(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1140 int test_buffer_write_uint( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1141 {
1142 cl_uint *inptr[5];
1143 size_t ptrSizes[5];
1144 int i, err;
1145 cl_uint j;
1146 MTdata d = init_genrand( gRandomSeed );
1147 int (*foo)(void *,void *,int);
1148
1149 size_t min_alignment = get_min_alignment(context);
1150
1151 foo = verify_write_uint;
1152
1153 ptrSizes[0] = sizeof(cl_uint);
1154 ptrSizes[1] = ptrSizes[0] << 1;
1155 ptrSizes[2] = ptrSizes[1] << 1;
1156 ptrSizes[3] = ptrSizes[2] << 1;
1157 ptrSizes[4] = ptrSizes[3] << 1;
1158
1159 for ( i = 0; i < 5; i++ ){
1160 inptr[i] = (cl_uint *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1161
1162 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1163 inptr[i][j] = genrand_int32(d);
1164 }
1165
1166 err = test_buffer_write( deviceID, context, queue, num_elements, sizeof( cl_uint ), (char*)"uint", 5, (void**)inptr,
1167 buffer_write_uint_kernel_code, uint_kernel_name, foo, d );
1168
1169 for ( i = 0; i < 5; i++ ){
1170 align_free( (void *)inptr[i] );
1171 }
1172
1173 free_mtdata(d);
1174 return err;
1175
1176 } // end test_buffer_uint_write()
1177
1178
test_buffer_write_short(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1179 int test_buffer_write_short( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1180 {
1181 short *inptr[5];
1182 size_t ptrSizes[5];
1183 int i, err;
1184 cl_uint j;
1185 MTdata d = init_genrand( gRandomSeed );
1186 int (*foo)(void *,void *,int);
1187
1188 size_t min_alignment = get_min_alignment(context);
1189
1190 foo = verify_write_short;
1191
1192 ptrSizes[0] = sizeof(cl_short);
1193 ptrSizes[1] = ptrSizes[0] << 1;
1194 ptrSizes[2] = ptrSizes[1] << 1;
1195 ptrSizes[3] = ptrSizes[2] << 1;
1196 ptrSizes[4] = ptrSizes[3] << 1;
1197
1198 for ( i = 0; i < 5; i++ ){
1199 inptr[i] = (cl_short *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1200
1201 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1202 inptr[i][j] = (cl_short)genrand_int32(d);
1203 }
1204
1205 err = test_buffer_write( deviceID, context, queue, num_elements, sizeof( cl_short ), (char*)"short", 5, (void**)inptr,
1206 buffer_write_short_kernel_code, short_kernel_name, foo, d );
1207
1208 for ( i = 0; i < 5; i++ ){
1209 align_free( (void *)inptr[i] );
1210
1211 }
1212
1213 free_mtdata(d);
1214 return err;
1215
1216 } // end test_buffer_short_write()
1217
1218
test_buffer_write_ushort(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1219 int test_buffer_write_ushort( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1220 {
1221 cl_ushort *inptr[5];
1222 size_t ptrSizes[5];
1223 int i, err;
1224 cl_uint j;
1225 MTdata d = init_genrand( gRandomSeed );
1226 int (*foo)(void *,void *,int);
1227
1228 size_t min_alignment = get_min_alignment(context);
1229
1230 foo = verify_write_ushort;
1231
1232 ptrSizes[0] = sizeof(cl_ushort);
1233 ptrSizes[1] = ptrSizes[0] << 1;
1234 ptrSizes[2] = ptrSizes[1] << 1;
1235 ptrSizes[3] = ptrSizes[2] << 1;
1236 ptrSizes[4] = ptrSizes[3] << 1;
1237
1238 for ( i = 0; i < 5; i++ ){
1239 inptr[i] = (cl_ushort *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1240
1241 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1242 inptr[i][j] = (cl_ushort)genrand_int32(d);
1243 }
1244
1245 err = test_buffer_write( deviceID, context, queue, num_elements, sizeof( cl_ushort ), (char*)"ushort", 5, (void**)inptr,
1246 buffer_write_ushort_kernel_code, ushort_kernel_name, foo, d );
1247
1248 for ( i = 0; i < 5; i++ ){
1249 align_free( (void *)inptr[i] );
1250
1251 }
1252
1253 free_mtdata(d);
1254 return err;
1255
1256 } // end test_buffer_ushort_write()
1257
1258
test_buffer_write_char(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1259 int test_buffer_write_char( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1260 {
1261 char *inptr[5];
1262 size_t ptrSizes[5];
1263 int i, err;
1264 cl_uint j;
1265 MTdata d = init_genrand( gRandomSeed );
1266 int (*foo)(void *,void *,int);
1267
1268 size_t min_alignment = get_min_alignment(context);
1269
1270 foo = verify_write_char;
1271
1272 ptrSizes[0] = sizeof(cl_char);
1273 ptrSizes[1] = ptrSizes[0] << 1;
1274 ptrSizes[2] = ptrSizes[1] << 1;
1275 ptrSizes[3] = ptrSizes[2] << 1;
1276 ptrSizes[4] = ptrSizes[3] << 1;
1277
1278 for ( i = 0; i < 5; i++ ){
1279 inptr[i] = (char *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1280
1281 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1282 inptr[i][j] = (char)genrand_int32(d);
1283 }
1284
1285 err = test_buffer_write( deviceID, context, queue, num_elements, sizeof( cl_char ), (char*)"char", 5, (void**)inptr,
1286 buffer_write_char_kernel_code, char_kernel_name, foo, d );
1287
1288 for ( i = 0; i < 5; i++ ){
1289 align_free( (void *)inptr[i] );
1290
1291 }
1292
1293 free_mtdata(d);
1294 return err;
1295
1296 } // end test_buffer_char_write()
1297
1298
test_buffer_write_uchar(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1299 int test_buffer_write_uchar( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1300 {
1301 uchar *inptr[5];
1302 size_t ptrSizes[5];
1303 int i, err;
1304 cl_uint j;
1305 MTdata d = init_genrand( gRandomSeed );
1306 int (*foo)(void *,void *,int);
1307
1308 size_t min_alignment = get_min_alignment(context);
1309
1310 foo = verify_write_uchar;
1311
1312 ptrSizes[0] = sizeof(cl_uchar);
1313 ptrSizes[1] = ptrSizes[0] << 1;
1314 ptrSizes[2] = ptrSizes[1] << 1;
1315 ptrSizes[3] = ptrSizes[2] << 1;
1316 ptrSizes[4] = ptrSizes[3] << 1;
1317
1318 for ( i = 0; i < 5; i++ ){
1319 inptr[i] = (uchar *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1320
1321 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1322 inptr[i][j] = (uchar)genrand_int32(d);
1323 }
1324
1325 err = test_buffer_write( deviceID, context, queue, num_elements, sizeof( cl_uchar ), (char*)"uchar", 5, (void**)inptr,
1326 buffer_write_uchar_kernel_code, uchar_kernel_name, foo, d );
1327
1328 for ( i = 0; i < 5; i++ ){
1329 align_free( (void *)inptr[i] );
1330
1331 }
1332
1333 free_mtdata(d);
1334 return err;
1335
1336 } // end test_buffer_uchar_write()
1337
1338
test_buffer_write_float(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1339 int test_buffer_write_float( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1340 {
1341 float *inptr[5];
1342 size_t ptrSizes[5];
1343 int i, err;
1344 cl_uint j;
1345 MTdata d = init_genrand( gRandomSeed );
1346 int (*foo)(void *,void *,int);
1347
1348 size_t min_alignment = get_min_alignment(context);
1349
1350 foo = verify_write_float;
1351
1352 ptrSizes[0] = sizeof(cl_float);
1353 ptrSizes[1] = ptrSizes[0] << 1;
1354 ptrSizes[2] = ptrSizes[1] << 1;
1355 ptrSizes[3] = ptrSizes[2] << 1;
1356 ptrSizes[4] = ptrSizes[3] << 1;
1357
1358 for ( i = 0; i < 5; i++ ){
1359 inptr[i] = (float *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1360
1361 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1362 inptr[i][j] = get_random_float( -FLT_MAX, FLT_MAX, d );
1363 }
1364
1365 err = test_buffer_write( deviceID, context, queue, num_elements, sizeof( cl_float ), (char*)"float", 5, (void**)inptr,
1366 buffer_write_float_kernel_code, float_kernel_name, foo, d );
1367
1368 for ( i = 0; i < 5; i++ ){
1369 align_free( (void *)inptr[i] );
1370 }
1371
1372 free_mtdata(d);
1373 return err;
1374
1375 } // end test_buffer_float_write()
1376
1377
test_buffer_write_half(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1378 int test_buffer_write_half( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1379 {
1380 PASSIVE_REQUIRE_FP16_SUPPORT(deviceID)
1381 float *inptr[5];
1382 size_t ptrSizes[5];
1383 int i, err;
1384 cl_uint j;
1385 MTdata d = init_genrand( gRandomSeed );
1386 int (*foo)(void *,void *,int);
1387
1388 size_t min_alignment = get_min_alignment(context);
1389
1390 foo = verify_write_half;
1391
1392 ptrSizes[0] = sizeof( cl_float ) / 2;
1393 ptrSizes[1] = ptrSizes[0] << 1;
1394 ptrSizes[2] = ptrSizes[1] << 1;
1395 ptrSizes[3] = ptrSizes[2] << 1;
1396 ptrSizes[4] = ptrSizes[3] << 1;
1397
1398 for ( i = 0; i < 5; i++ ){
1399 inptr[i] = (float *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1400
1401 for ( j = 0; j < ptrSizes[i] * num_elements / ( ptrSizes[0] * 2 ); j++ )
1402 inptr[i][j] = get_random_float( -FLT_MAX, FLT_MAX, d );
1403 }
1404
1405 err = test_buffer_write(deviceID, context, queue, num_elements,
1406 sizeof(cl_half), (char *)"half", 5, (void **)inptr,
1407 buffer_write_half_kernel_code, half_kernel_name,
1408 foo, d);
1409
1410 for ( i = 0; i < 5; i++ ){
1411 align_free( (void *)inptr[i] );
1412 }
1413
1414 free_mtdata(d);
1415 return err;
1416
1417 } // end test_buffer_half_write()
1418
1419
test_buffer_write_long(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1420 int test_buffer_write_long( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1421 {
1422 cl_long *inptr[5];
1423 size_t ptrSizes[5];
1424 int i, err;
1425 cl_uint j;
1426 MTdata d = init_genrand( gRandomSeed );
1427 int (*foo)(void *,void *,int);
1428
1429 size_t min_alignment = get_min_alignment(context);
1430
1431 foo = verify_write_long;
1432
1433 ptrSizes[0] = sizeof(cl_long);
1434 ptrSizes[1] = ptrSizes[0] << 1;
1435 ptrSizes[2] = ptrSizes[1] << 1;
1436 ptrSizes[3] = ptrSizes[2] << 1;
1437 ptrSizes[4] = ptrSizes[3] << 1;
1438
1439 //skip devices that don't support long
1440 if (! gHasLong )
1441 {
1442 log_info( "Device does not support 64-bit integers. Skipping test.\n" );
1443 return CL_SUCCESS;
1444 }
1445
1446 for ( i = 0; i < 5; i++ ){
1447 inptr[i] = (cl_long *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1448
1449 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1450 inptr[i][j] = (cl_long) genrand_int32(d) ^ ((cl_long) genrand_int32(d) << 32);
1451 }
1452
1453 err = test_buffer_write( deviceID, context, queue, num_elements, sizeof( cl_long ), (char*)"cl_long", 5, (void**)inptr,
1454 buffer_write_long_kernel_code, long_kernel_name, foo, d );
1455
1456 for ( i = 0; i < 5; i++ ){
1457 align_free( (void *)inptr[i] );
1458 }
1459
1460 free_mtdata(d);
1461 return err;
1462
1463 } // end test_buffer_long_write()
1464
1465
test_buffer_write_ulong(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1466 int test_buffer_write_ulong( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1467 {
1468 cl_ulong *inptr[5];
1469 size_t ptrSizes[5];
1470 int i, err;
1471 cl_uint j;
1472 MTdata d = init_genrand( gRandomSeed );
1473 int (*foo)(void *,void *,int);
1474
1475 size_t min_alignment = get_min_alignment(context);
1476
1477 foo = verify_write_ulong;
1478
1479 ptrSizes[0] = sizeof(cl_ulong);
1480 ptrSizes[1] = ptrSizes[0] << 1;
1481 ptrSizes[2] = ptrSizes[1] << 1;
1482 ptrSizes[3] = ptrSizes[2] << 1;
1483 ptrSizes[4] = ptrSizes[3] << 1;
1484
1485 if (! gHasLong )
1486 {
1487 log_info( "Device does not support 64-bit integers. Skipping test.\n" );
1488 return CL_SUCCESS;
1489 }
1490
1491 for ( i = 0; i < 5; i++ ){
1492 inptr[i] = (cl_ulong *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1493
1494 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1495 inptr[i][j] = (cl_ulong) genrand_int32(d) | ((cl_ulong) genrand_int32(d) << 32);
1496 }
1497
1498 err = test_buffer_write( deviceID, context, queue, num_elements, sizeof( cl_ulong ), (char*)"ulong long", 5, (void**)inptr,
1499 buffer_write_ulong_kernel_code, ulong_kernel_name, foo, d );
1500
1501 for ( i = 0; i < 5; i++ ){
1502 align_free( (void *)inptr[i] );
1503 }
1504
1505 free_mtdata(d);
1506
1507 return err;
1508
1509 } // end test_buffer_ulong_write()
1510
1511
test_buffer_map_write_int(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1512 int test_buffer_map_write_int( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1513 {
1514 gTestMap = 1;
1515 return test_buffer_write_int(deviceID, context, queue, num_elements);
1516 }
1517
test_buffer_map_write_uint(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1518 int test_buffer_map_write_uint( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1519 {
1520 gTestMap = 1;
1521 return test_buffer_write_uint(deviceID, context, queue, num_elements);
1522 }
1523
test_buffer_map_write_long(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1524 int test_buffer_map_write_long( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1525 {
1526 gTestMap = 1;
1527 return test_buffer_write_long(deviceID, context, queue, num_elements);
1528 }
1529
test_buffer_map_write_ulong(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1530 int test_buffer_map_write_ulong( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1531 {
1532 gTestMap = 1;
1533 return test_buffer_write_ulong(deviceID, context, queue, num_elements);
1534 }
1535
test_buffer_map_write_short(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1536 int test_buffer_map_write_short( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1537 {
1538 gTestMap = 1;
1539 return test_buffer_write_short(deviceID, context, queue, num_elements);
1540 }
1541
test_buffer_map_write_ushort(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1542 int test_buffer_map_write_ushort( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1543 {
1544 gTestMap = 1;
1545 return test_buffer_write_ushort(deviceID, context, queue, num_elements);
1546 }
1547
test_buffer_map_write_char(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1548 int test_buffer_map_write_char( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1549 {
1550 gTestMap = 1;
1551 return test_buffer_write_char(deviceID, context, queue, num_elements);
1552 }
1553
test_buffer_map_write_uchar(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1554 int test_buffer_map_write_uchar( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1555 {
1556 gTestMap = 1;
1557 return test_buffer_write_uchar(deviceID, context, queue, num_elements);
1558 }
1559
test_buffer_map_write_float(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1560 int test_buffer_map_write_float( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1561 {
1562 gTestMap = 1;
1563 return test_buffer_write_float(deviceID, context, queue, num_elements);
1564 }
1565
test_buffer_map_write_struct(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1566 int test_buffer_map_write_struct( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1567 {
1568 gTestMap = 1;
1569 return test_buffer_write_struct(deviceID, context, queue, num_elements);
1570 }
1571
1572
test_buffer_write_async_int(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1573 int test_buffer_write_async_int( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1574 {
1575 int *inptr[5];
1576 size_t ptrSizes[5];
1577 int i, err;
1578 cl_uint j;
1579 MTdata d = init_genrand( gRandomSeed );
1580 int (*foo)(void *,void *,int);
1581
1582 size_t min_alignment = get_min_alignment(context);
1583
1584 foo = verify_write_int;
1585
1586 ptrSizes[0] = sizeof(cl_int);
1587 ptrSizes[1] = ptrSizes[0] << 1;
1588 ptrSizes[2] = ptrSizes[1] << 1;
1589 ptrSizes[3] = ptrSizes[2] << 1;
1590 ptrSizes[4] = ptrSizes[3] << 1;
1591
1592 for ( i = 0; i < 5; i++ ){
1593 inptr[i] = (int *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1594
1595 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1596 inptr[i][j] = (int)genrand_int32(d);
1597 }
1598
1599 err = test_buffer_write_array_async( deviceID, context, queue, num_elements, sizeof( cl_int ), (char*)"int", 5, (void**)inptr,
1600 buffer_write_int_kernel_code, int_kernel_name, foo );
1601
1602 for ( i = 0; i < 5; i++ ){
1603 align_free( (void *)inptr[i] );
1604 }
1605
1606 free_mtdata(d);
1607 return err;
1608
1609 } // end test_buffer_int_write_array_async()
1610
1611
test_buffer_write_async_uint(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1612 int test_buffer_write_async_uint( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1613 {
1614 cl_uint *inptr[5];
1615 size_t ptrSizes[5];
1616 int i, err;
1617 cl_uint j;
1618 MTdata d = init_genrand( gRandomSeed );
1619 int (*foo)(void *,void *,int);
1620
1621 size_t min_alignment = get_min_alignment(context);
1622
1623 foo = verify_write_uint;
1624
1625 ptrSizes[0] = sizeof(cl_uint);
1626 ptrSizes[1] = ptrSizes[0] << 1;
1627 ptrSizes[2] = ptrSizes[1] << 1;
1628 ptrSizes[3] = ptrSizes[2] << 1;
1629 ptrSizes[4] = ptrSizes[3] << 1;
1630
1631 for ( i = 0; i < 5; i++ ){
1632 inptr[i] = (cl_uint *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1633
1634 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1635 inptr[i][j] = (cl_uint)genrand_int32(d);
1636 }
1637
1638 err = test_buffer_write_array_async( deviceID, context, queue, num_elements, sizeof( cl_uint ), (char*)"uint", 5, (void**)inptr,
1639 buffer_write_uint_kernel_code, uint_kernel_name, foo );
1640
1641 for ( i = 0; i < 5; i++ ){
1642 align_free( (void *)inptr[i] );
1643 }
1644
1645 free_mtdata(d);
1646 return err;
1647
1648 } // end test_buffer_uint_write_array_async()
1649
1650
test_buffer_write_async_short(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1651 int test_buffer_write_async_short( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1652 {
1653 short *inptr[5];
1654 size_t ptrSizes[5];
1655 int i, err;
1656 cl_uint j;
1657 MTdata d = init_genrand( gRandomSeed );
1658 int (*foo)(void *,void *,int);
1659
1660 size_t min_alignment = get_min_alignment(context);
1661
1662 foo = verify_write_short;
1663
1664 ptrSizes[0] = sizeof(cl_short);
1665 ptrSizes[1] = ptrSizes[0] << 1;
1666 ptrSizes[2] = ptrSizes[1] << 1;
1667 ptrSizes[3] = ptrSizes[2] << 1;
1668 ptrSizes[4] = ptrSizes[3] << 1;
1669
1670 for ( i = 0; i < 5; i++ ){
1671 inptr[i] = (short *)align_malloc(ptrSizes[i] * num_elements + min_alignment, min_alignment);
1672
1673 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1674 inptr[i][j] = (short)genrand_int32(d);
1675 }
1676
1677 err = test_buffer_write_array_async( deviceID, context, queue, num_elements, sizeof( cl_short ), (char*)"short", 5, (void**)inptr,
1678 buffer_write_short_kernel_code, short_kernel_name, foo );
1679
1680 for ( i = 0; i < 5; i++ ){
1681 align_free( (void *)inptr[i] );
1682
1683 }
1684
1685 free_mtdata(d);
1686 return err;
1687
1688 } // end test_buffer_short_write_array_async()
1689
1690
test_buffer_write_async_ushort(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1691 int test_buffer_write_async_ushort( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1692 {
1693 cl_ushort *inptr[5];
1694 size_t ptrSizes[5];
1695 int i, err;
1696 cl_uint j;
1697 MTdata d = init_genrand( gRandomSeed );
1698 int (*foo)(void *,void *,int);
1699
1700 size_t min_alignment = get_min_alignment(context);
1701
1702 foo = verify_write_ushort;
1703
1704 ptrSizes[0] = sizeof(cl_ushort);
1705 ptrSizes[1] = ptrSizes[0] << 1;
1706 ptrSizes[2] = ptrSizes[1] << 1;
1707 ptrSizes[3] = ptrSizes[2] << 1;
1708 ptrSizes[4] = ptrSizes[3] << 1;
1709
1710 for ( i = 0; i < 5; i++ ){
1711 inptr[i] = (cl_ushort *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1712
1713 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1714 inptr[i][j] = (cl_ushort)genrand_int32(d);
1715 }
1716
1717 err = test_buffer_write_array_async( deviceID, context, queue, num_elements, sizeof( cl_ushort ), (char*)"ushort", 5, (void**)inptr,
1718 buffer_write_ushort_kernel_code, ushort_kernel_name, foo );
1719
1720 for ( i = 0; i < 5; i++ ){
1721 align_free( (void *)inptr[i] );
1722
1723 }
1724
1725 free_mtdata(d);
1726 return err;
1727
1728 } // end test_buffer_ushort_write_array_async()
1729
1730
test_buffer_write_async_char(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1731 int test_buffer_write_async_char( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1732 {
1733 char *inptr[5];
1734 size_t ptrSizes[5];
1735 int i, err;
1736 cl_uint j;
1737 MTdata d = init_genrand( gRandomSeed );
1738 int (*foo)(void *,void *,int);
1739
1740 size_t min_alignment = get_min_alignment(context);
1741
1742 foo = verify_write_char;
1743
1744 ptrSizes[0] = sizeof(cl_char);
1745 ptrSizes[1] = ptrSizes[0] << 1;
1746 ptrSizes[2] = ptrSizes[1] << 1;
1747 ptrSizes[3] = ptrSizes[2] << 1;
1748 ptrSizes[4] = ptrSizes[3] << 1;
1749
1750 for ( i = 0; i < 5; i++ ){
1751 inptr[i] = (char *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1752
1753 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1754 inptr[i][j] = (char)genrand_int32(d);
1755 }
1756
1757 err = test_buffer_write_array_async( deviceID, context, queue, num_elements, sizeof( cl_char ), (char*)"char", 5, (void**)inptr,
1758 buffer_write_char_kernel_code, char_kernel_name, foo );
1759
1760 for ( i = 0; i < 5; i++ ){
1761 align_free( (void *)inptr[i] );
1762
1763 }
1764
1765 free_mtdata(d);
1766 return err;
1767
1768 } // end test_buffer_char_write_array_async()
1769
1770
test_buffer_write_async_uchar(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1771 int test_buffer_write_async_uchar( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1772 {
1773 uchar *inptr[5];
1774 size_t ptrSizes[5];
1775 int i, err;
1776 cl_uint j;
1777 MTdata d = init_genrand( gRandomSeed );
1778 int (*foo)(void *,void *,int);
1779
1780 size_t min_alignment = get_min_alignment(context);
1781
1782 foo = verify_write_uchar;
1783
1784 ptrSizes[0] = sizeof(cl_uchar);
1785 ptrSizes[1] = ptrSizes[0] << 1;
1786 ptrSizes[2] = ptrSizes[1] << 1;
1787 ptrSizes[3] = ptrSizes[2] << 1;
1788 ptrSizes[4] = ptrSizes[3] << 1;
1789
1790 for ( i = 0; i < 5; i++ ){
1791 inptr[i] = (uchar *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1792
1793 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1794 inptr[i][j] = (uchar)genrand_int32(d);
1795 }
1796
1797 err = test_buffer_write_array_async( deviceID, context, queue, num_elements, sizeof( cl_uchar ), (char*)"uchar", 5, (void**)inptr,
1798 buffer_write_uchar_kernel_code, uchar_kernel_name, foo );
1799
1800 for ( i = 0; i < 5; i++ ){
1801 align_free( (void *)inptr[i] );
1802
1803 }
1804
1805 free_mtdata(d);
1806 return err;
1807
1808 } // end test_buffer_uchar_write_array_async()
1809
1810
test_buffer_write_async_float(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1811 int test_buffer_write_async_float( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1812 {
1813 float *inptr[5];
1814 size_t ptrSizes[5];
1815 int i, err;
1816 cl_uint j;
1817 MTdata d = init_genrand( gRandomSeed );
1818 int (*foo)(void *,void *,int);
1819
1820 size_t min_alignment = get_min_alignment(context);
1821
1822 foo = verify_write_float;
1823
1824 ptrSizes[0] = sizeof(cl_float);
1825 ptrSizes[1] = ptrSizes[0] << 1;
1826 ptrSizes[2] = ptrSizes[1] << 1;
1827 ptrSizes[3] = ptrSizes[2] << 1;
1828 ptrSizes[4] = ptrSizes[3] << 1;
1829
1830 for ( i = 0; i < 5; i++ ){
1831 inptr[i] = (float *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1832
1833 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1834 inptr[i][j] = get_random_float( -FLT_MAX, FLT_MAX, d );
1835 }
1836
1837 err = test_buffer_write_array_async( deviceID, context, queue, num_elements, sizeof( cl_float ), (char*)"float", 5, (void**)inptr,
1838 buffer_write_float_kernel_code, float_kernel_name, foo );
1839
1840 for ( i = 0; i < 5; i++ ){
1841 align_free( (void *)inptr[i] );
1842 }
1843
1844 free_mtdata(d);
1845 return err;
1846
1847 } // end test_buffer_float_write_array_async()
1848
1849
test_buffer_write_async_long(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1850 int test_buffer_write_async_long( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1851 {
1852 cl_long *inptr[5];
1853 size_t ptrSizes[5];
1854 int i, err;
1855 cl_uint j;
1856 MTdata d = init_genrand( gRandomSeed );
1857 int (*foo)(void *,void *,int);
1858
1859 size_t min_alignment = get_min_alignment(context);
1860
1861 foo = verify_write_long;
1862
1863 ptrSizes[0] = sizeof(cl_long);
1864 ptrSizes[1] = ptrSizes[0] << 1;
1865 ptrSizes[2] = ptrSizes[1] << 1;
1866 ptrSizes[3] = ptrSizes[2] << 1;
1867 ptrSizes[4] = ptrSizes[3] << 1;
1868
1869 if (! gHasLong )
1870 {
1871 log_info( "Device does not support 64-bit integers. Skipping test.\n" );
1872 return CL_SUCCESS;
1873 }
1874
1875 for ( i = 0; i < 5; i++ ){
1876 inptr[i] = (cl_long *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1877
1878 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1879 inptr[i][j] = ((cl_long) genrand_int32(d)) ^ ((cl_long) genrand_int32(d) << 32);
1880 }
1881
1882 err = test_buffer_write_array_async( deviceID, context, queue, num_elements, sizeof( cl_long ), (char*)"cl_long", 5, (void**)inptr,
1883 buffer_write_long_kernel_code, long_kernel_name, foo );
1884
1885 for ( i = 0; i < 5; i++ ){
1886 align_free( (void *)inptr[i] );
1887 }
1888
1889 free_mtdata(d);
1890 return err;
1891
1892 } // end test_buffer_long_write_array_async()
1893
1894
test_buffer_write_async_ulong(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1895 int test_buffer_write_async_ulong( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
1896 {
1897 cl_ulong *inptr[5];
1898 size_t ptrSizes[5];
1899 int i, err;
1900 cl_uint j;
1901 MTdata d = init_genrand( gRandomSeed );
1902 int (*foo)(void *,void *,int);
1903
1904 size_t min_alignment = get_min_alignment(context);
1905
1906 foo = verify_write_ulong;
1907
1908 ptrSizes[0] = sizeof(cl_ulong);
1909 ptrSizes[1] = ptrSizes[0] << 1;
1910 ptrSizes[2] = ptrSizes[1] << 1;
1911 ptrSizes[3] = ptrSizes[2] << 1;
1912 ptrSizes[4] = ptrSizes[3] << 1;
1913
1914 if (! gHasLong )
1915 {
1916 log_info( "Device does not support 64-bit integers. Skipping test.\n" );
1917 return CL_SUCCESS;
1918 }
1919
1920 for ( i = 0; i < 5; i++ ){
1921 inptr[i] = (cl_ulong *)align_malloc(ptrSizes[i] * num_elements, min_alignment);
1922
1923 for ( j = 0; j < ptrSizes[i] * num_elements / ptrSizes[0]; j++ )
1924 inptr[i][j] = (cl_ulong) genrand_int32(d) | ((cl_ulong) genrand_int32(d) << 32);
1925 }
1926
1927 err = test_buffer_write_array_async( deviceID, context, queue, num_elements, sizeof( cl_ulong ), (char*)"ulong long", 5, (void**)inptr,
1928 buffer_write_ulong_kernel_code, ulong_kernel_name, foo );
1929
1930 for ( i = 0; i < 5; i++ ){
1931 align_free( (void *)inptr[i] );
1932 }
1933
1934 free_mtdata(d);
1935 return err;
1936
1937 } // end test_buffer_ulong_write_array_async()
1938
1939