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 <string.h>
20 #include <time.h>
21 #include <sys/types.h>
22 #include <sys/stat.h>
23
24 #include "procs.h"
25 #include "harness/testHarness.h"
26
27 #define TEST_PRIME_INT ((1<<16)+1)
28 #define TEST_PRIME_UINT ((1U<<16)+1U)
29 #define TEST_PRIME_LONG ((1LL<<32)+1LL)
30 #define TEST_PRIME_ULONG ((1ULL<<32)+1ULL)
31 #define TEST_PRIME_SHORT ((1S<<8)+1S)
32 #define TEST_PRIME_FLOAT (float)3.40282346638528860e+38
33 #define TEST_PRIME_HALF 119.f
34 #define TEST_BOOL true
35 #define TEST_PRIME_CHAR 0x77
36
37
38 #ifndef ulong
39 typedef unsigned long ulong;
40 #endif
41
42 #ifndef uchar
43 typedef unsigned char uchar;
44 #endif
45
46 #ifndef TestStruct
47 typedef struct{
48 int a;
49 float b;
50 } TestStruct;
51 #endif
52
53
54
55 //--- the code for the kernel executables
56 static const char *stream_read_int_kernel_code[] = {
57 "__kernel void test_stream_read_int(__global int *dst)\n"
58 "{\n"
59 " int tid = get_global_id(0);\n"
60 "\n"
61 " dst[tid] = ((1<<16)+1);\n"
62 "}\n",
63
64 "__kernel void test_stream_read_int2(__global int2 *dst)\n"
65 "{\n"
66 " int tid = get_global_id(0);\n"
67 "\n"
68 " dst[tid] = ((1<<16)+1);\n"
69 "}\n",
70
71 "__kernel void test_stream_read_int4(__global int4 *dst)\n"
72 "{\n"
73 " int tid = get_global_id(0);\n"
74 "\n"
75 " dst[tid] = ((1<<16)+1);\n"
76 "}\n",
77
78 "__kernel void test_stream_read_int8(__global int8 *dst)\n"
79 "{\n"
80 " int tid = get_global_id(0);\n"
81 "\n"
82 " dst[tid] = ((1<<16)+1);\n"
83 "}\n",
84
85 "__kernel void test_stream_read_int16(__global int16 *dst)\n"
86 "{\n"
87 " int tid = get_global_id(0);\n"
88 "\n"
89 " dst[tid] = ((1<<16)+1);\n"
90 "}\n" };
91
92 static const char *int_kernel_name[] = { "test_stream_read_int", "test_stream_read_int2", "test_stream_read_int4", "test_stream_read_int8", "test_stream_read_int16" };
93
94 const char *stream_read_uint_kernel_code[] = {
95 "__kernel void test_stream_read_uint(__global uint *dst)\n"
96 "{\n"
97 " int tid = get_global_id(0);\n"
98 "\n"
99 " dst[tid] = ((1U<<16)+1U);\n"
100 "}\n",
101
102 "__kernel void test_stream_read_uint2(__global uint2 *dst)\n"
103 "{\n"
104 " int tid = get_global_id(0);\n"
105 "\n"
106 " dst[tid] = ((1U<<16)+1U);\n"
107 "}\n",
108
109 "__kernel void test_stream_read_uint4(__global uint4 *dst)\n"
110 "{\n"
111 " int tid = get_global_id(0);\n"
112 "\n"
113 " dst[tid] = ((1U<<16)+1U);\n"
114 "}\n",
115
116 "__kernel void test_stream_read_uint8(__global uint8 *dst)\n"
117 "{\n"
118 " int tid = get_global_id(0);\n"
119 "\n"
120 " dst[tid] = ((1U<<16)+1U);\n"
121 "}\n",
122
123 "__kernel void test_stream_read_uint16(__global uint16 *dst)\n"
124 "{\n"
125 " int tid = get_global_id(0);\n"
126 "\n"
127 " dst[tid] = ((1U<<16)+1U);\n"
128 "}\n" };
129
130 const char *uint_kernel_name[] = { "test_stream_read_uint", "test_stream_read_uint2", "test_stream_read_uint4", "test_stream_read_uint8", "test_stream_read_uint16" };
131
132 const char *stream_read_long_kernel_code[] = {
133 "__kernel void test_stream_read_long(__global long *dst)\n"
134 "{\n"
135 " int tid = get_global_id(0);\n"
136 "\n"
137 " dst[tid] = ((1L<<32)+1L);\n"
138 "}\n",
139
140 "__kernel void test_stream_read_long2(__global long2 *dst)\n"
141 "{\n"
142 " int tid = get_global_id(0);\n"
143 "\n"
144 " dst[tid] = ((1L<<32)+1L);\n"
145 "}\n",
146
147 "__kernel void test_stream_read_long4(__global long4 *dst)\n"
148 "{\n"
149 " int tid = get_global_id(0);\n"
150 "\n"
151 " dst[tid] = ((1L<<32)+1L);\n"
152 "}\n",
153
154 "__kernel void test_stream_read_long8(__global long8 *dst)\n"
155 "{\n"
156 " int tid = get_global_id(0);\n"
157 "\n"
158 " dst[tid] = ((1L<<32)+1L);\n"
159 "}\n",
160
161 "__kernel void test_stream_read_long16(__global long16 *dst)\n"
162 "{\n"
163 " int tid = get_global_id(0);\n"
164 "\n"
165 " dst[tid] = ((1L<<32)+1L);\n"
166 "}\n" };
167
168 const char *long_kernel_name[] = { "test_stream_read_long", "test_stream_read_long2", "test_stream_read_long4", "test_stream_read_long8", "test_stream_read_long16" };
169
170 const char *stream_read_ulong_kernel_code[] = {
171 "__kernel void test_stream_read_ulong(__global ulong *dst)\n"
172 "{\n"
173 " int tid = get_global_id(0);\n"
174 "\n"
175 " dst[tid] = ((1UL<<32)+1UL);\n"
176 "}\n",
177
178 "__kernel void test_stream_read_ulong2(__global ulong2 *dst)\n"
179 "{\n"
180 " int tid = get_global_id(0);\n"
181 "\n"
182 " dst[tid] = ((1UL<<32)+1UL);\n"
183 "}\n",
184
185 "__kernel void test_stream_read_ulong4(__global ulong4 *dst)\n"
186 "{\n"
187 " int tid = get_global_id(0);\n"
188 "\n"
189 " dst[tid] = ((1UL<<32)+1UL);\n"
190 "}\n",
191
192 "__kernel void test_stream_read_ulong8(__global ulong8 *dst)\n"
193 "{\n"
194 " int tid = get_global_id(0);\n"
195 "\n"
196 " dst[tid] = ((1UL<<32)+1UL);\n"
197 "}\n",
198
199 "__kernel void test_stream_read_ulong16(__global ulong16 *dst)\n"
200 "{\n"
201 " int tid = get_global_id(0);\n"
202 "\n"
203 " dst[tid] = ((1UL<<32)+1UL);\n"
204 "}\n" };
205
206 const char *ulong_kernel_name[] = { "test_stream_read_ulong", "test_stream_read_ulong2", "test_stream_read_ulong4", "test_stream_read_ulong8", "test_stream_read_ulong16" };
207
208 const char *stream_read_short_kernel_code[] = {
209 "__kernel void test_stream_read_short(__global short *dst)\n"
210 "{\n"
211 " int tid = get_global_id(0);\n"
212 "\n"
213 " dst[tid] = (short)((1<<8)+1);\n"
214 "}\n",
215
216 "__kernel void test_stream_read_short2(__global short2 *dst)\n"
217 "{\n"
218 " int tid = get_global_id(0);\n"
219 "\n"
220 " dst[tid] = (short)((1<<8)+1);\n"
221 "}\n",
222
223 "__kernel void test_stream_read_short4(__global short4 *dst)\n"
224 "{\n"
225 " int tid = get_global_id(0);\n"
226 "\n"
227 " dst[tid] = (short)((1<<8)+1);\n"
228 "}\n",
229
230 "__kernel void test_stream_read_short8(__global short8 *dst)\n"
231 "{\n"
232 " int tid = get_global_id(0);\n"
233 "\n"
234 " dst[tid] = (short)((1<<8)+1);\n"
235 "}\n",
236
237 "__kernel void test_stream_read_short16(__global short16 *dst)\n"
238 "{\n"
239 " int tid = get_global_id(0);\n"
240 "\n"
241 " dst[tid] = (short)((1<<8)+1);\n"
242 "}\n" };
243
244 const char *short_kernel_name[] = { "test_stream_read_short", "test_stream_read_short2", "test_stream_read_short4", "test_stream_read_short8", "test_stream_read_short16" };
245
246
247 const char *stream_read_ushort_kernel_code[] = {
248 "__kernel void test_stream_read_ushort(__global ushort *dst)\n"
249 "{\n"
250 " int tid = get_global_id(0);\n"
251 "\n"
252 " dst[tid] = (ushort)((1<<8)+1);\n"
253 "}\n",
254
255 "__kernel void test_stream_read_ushort2(__global ushort2 *dst)\n"
256 "{\n"
257 " int tid = get_global_id(0);\n"
258 "\n"
259 " dst[tid] = (ushort)((1<<8)+1);\n"
260 "}\n",
261
262 "__kernel void test_stream_read_ushort4(__global ushort4 *dst)\n"
263 "{\n"
264 " int tid = get_global_id(0);\n"
265 "\n"
266 " dst[tid] = (ushort)((1<<8)+1);\n"
267 "}\n",
268
269 "__kernel void test_stream_read_ushort8(__global ushort8 *dst)\n"
270 "{\n"
271 " int tid = get_global_id(0);\n"
272 "\n"
273 " dst[tid] = (ushort)((1<<8)+1);\n"
274 "}\n",
275
276 "__kernel void test_stream_read_ushort16(__global ushort16 *dst)\n"
277 "{\n"
278 " int tid = get_global_id(0);\n"
279 "\n"
280 " dst[tid] = (ushort)((1<<8)+1);\n"
281 "}\n" };
282
283 static const char *ushort_kernel_name[] = { "test_stream_read_ushort", "test_stream_read_ushort2", "test_stream_read_ushort4", "test_stream_read_ushort8", "test_stream_read_ushort16" };
284
285
286 const char *stream_read_float_kernel_code[] = {
287 "__kernel void test_stream_read_float(__global float *dst)\n"
288 "{\n"
289 " int tid = get_global_id(0);\n"
290 "\n"
291 " dst[tid] = (float)3.40282346638528860e+38;\n"
292 "}\n",
293
294 "__kernel void test_stream_read_float2(__global float2 *dst)\n"
295 "{\n"
296 " int tid = get_global_id(0);\n"
297 "\n"
298 " dst[tid] = (float)3.40282346638528860e+38;\n"
299 "}\n",
300
301 "__kernel void test_stream_read_float4(__global float4 *dst)\n"
302 "{\n"
303 " int tid = get_global_id(0);\n"
304 "\n"
305 " dst[tid] = (float)3.40282346638528860e+38;\n"
306 "}\n",
307
308 "__kernel void test_stream_read_float8(__global float8 *dst)\n"
309 "{\n"
310 " int tid = get_global_id(0);\n"
311 "\n"
312 " dst[tid] = (float)3.40282346638528860e+38;\n"
313 "}\n",
314
315 "__kernel void test_stream_read_float16(__global float16 *dst)\n"
316 "{\n"
317 " int tid = get_global_id(0);\n"
318 "\n"
319 " dst[tid] = (float)3.40282346638528860e+38;\n"
320 "}\n" };
321
322 const char *float_kernel_name[] = { "test_stream_read_float", "test_stream_read_float2", "test_stream_read_float4", "test_stream_read_float8", "test_stream_read_float16" };
323
324
325 const char *stream_read_half_kernel_code[] = {
326 "__kernel void test_stream_read_half(__global half *dst)\n"
327 "{\n"
328 " int tid = get_global_id(0);\n"
329 "\n"
330 " dst[tid] = (half)119;\n"
331 "}\n",
332
333 "__kernel void test_stream_read_half2(__global half2 *dst)\n"
334 "{\n"
335 " int tid = get_global_id(0);\n"
336 "\n"
337 " dst[tid] = (half)119;\n"
338 "}\n",
339
340 "__kernel void test_stream_read_half4(__global half4 *dst)\n"
341 "{\n"
342 " int tid = get_global_id(0);\n"
343 "\n"
344 " dst[tid] = (half)119;\n"
345 "}\n",
346
347 "__kernel void test_stream_read_half8(__global half8 *dst)\n"
348 "{\n"
349 " int tid = get_global_id(0);\n"
350 "\n"
351 " dst[tid] = (half)119;\n"
352 "}\n",
353
354 "__kernel void test_stream_read_half16(__global half16 *dst)\n"
355 "{\n"
356 " int tid = get_global_id(0);\n"
357 "\n"
358 " dst[tid] = (half)119;\n"
359 "}\n" };
360
361 const char *half_kernel_name[] = { "test_stream_read_half", "test_stream_read_half2", "test_stream_read_half4", "test_stream_read_half8", "test_stream_read_half16" };
362
363
364 const char *stream_read_char_kernel_code[] = {
365 "__kernel void test_stream_read_char(__global char *dst)\n"
366 "{\n"
367 " int tid = get_global_id(0);\n"
368 "\n"
369 " dst[tid] = (char)'w';\n"
370 "}\n",
371
372 "__kernel void test_stream_read_char2(__global char2 *dst)\n"
373 "{\n"
374 " int tid = get_global_id(0);\n"
375 "\n"
376 " dst[tid] = (char)'w';\n"
377 "}\n",
378
379 "__kernel void test_stream_read_char4(__global char4 *dst)\n"
380 "{\n"
381 " int tid = get_global_id(0);\n"
382 "\n"
383 " dst[tid] = (char)'w';\n"
384 "}\n",
385
386 "__kernel void test_stream_read_char8(__global char8 *dst)\n"
387 "{\n"
388 " int tid = get_global_id(0);\n"
389 "\n"
390 " dst[tid] = (char)'w';\n"
391 "}\n",
392
393 "__kernel void test_stream_read_char16(__global char16 *dst)\n"
394 "{\n"
395 " int tid = get_global_id(0);\n"
396 "\n"
397 " dst[tid] = (char)'w';\n"
398 "}\n" };
399
400 const char *char_kernel_name[] = { "test_stream_read_char", "test_stream_read_char2", "test_stream_read_char4", "test_stream_read_char8", "test_stream_read_char16" };
401
402
403 const char *stream_read_uchar_kernel_code[] = {
404 "__kernel void test_stream_read_uchar(__global uchar *dst)\n"
405 "{\n"
406 " int tid = get_global_id(0);\n"
407 "\n"
408 " dst[tid] = 'w';\n"
409 "}\n",
410
411 "__kernel void test_stream_read_uchar2(__global uchar2 *dst)\n"
412 "{\n"
413 " int tid = get_global_id(0);\n"
414 "\n"
415 " dst[tid] = (uchar)'w';\n"
416 "}\n",
417
418 "__kernel void test_stream_read_uchar4(__global uchar4 *dst)\n"
419 "{\n"
420 " int tid = get_global_id(0);\n"
421 "\n"
422 " dst[tid] = (uchar)'w';\n"
423 "}\n",
424
425 "__kernel void test_stream_read_uchar8(__global uchar8 *dst)\n"
426 "{\n"
427 " int tid = get_global_id(0);\n"
428 "\n"
429 " dst[tid] = (uchar)'w';\n"
430 "}\n",
431
432 "__kernel void test_stream_read_uchar16(__global uchar16 *dst)\n"
433 "{\n"
434 " int tid = get_global_id(0);\n"
435 "\n"
436 " dst[tid] = (uchar)'w';\n"
437 "}\n" };
438
439 const char *uchar_kernel_name[] = { "test_stream_read_uchar", "test_stream_read_uchar2", "test_stream_read_uchar4", "test_stream_read_uchar8", "test_stream_read_uchar16" };
440
441
442 const char *stream_read_struct_kernel_code[] = {
443 "typedef struct{\n"
444 "int a;\n"
445 "float b;\n"
446 "} TestStruct;\n"
447 "__kernel void test_stream_read_struct(__global TestStruct *dst)\n"
448 "{\n"
449 " int tid = get_global_id(0);\n"
450 "\n"
451 " dst[tid].a = ((1<<16)+1);\n"
452 " dst[tid].b = (float)3.40282346638528860e+38;\n"
453 "}\n" };
454
455 const char *struct_kernel_name[] = { "test_stream_read_struct" };
456
457
458
459 //--- the verify functions
verify_read_int(void * ptr,int n)460 static int verify_read_int(void *ptr, int n)
461 {
462 int i;
463 int *outptr = (int *)ptr;
464
465 for (i=0; i<n; i++){
466 if( outptr[i] != TEST_PRIME_INT )
467 return -1;
468 }
469
470 return 0;
471 }
472
473
verify_read_uint(void * ptr,int n)474 static int verify_read_uint(void *ptr, int n)
475 {
476 int i;
477 cl_uint *outptr = (cl_uint *)ptr;
478
479 for (i=0; i<n; i++){
480 if( outptr[i] != TEST_PRIME_UINT )
481 return -1;
482 }
483
484 return 0;
485 }
486
487
verify_read_long(void * ptr,int n)488 static int verify_read_long(void *ptr, int n)
489 {
490 int i;
491 cl_long *outptr = (cl_long *)ptr;
492
493 for (i=0; i<n; i++){
494 if( outptr[i] != TEST_PRIME_LONG )
495 return -1;
496 }
497
498 return 0;
499 }
500
501
verify_read_ulong(void * ptr,int n)502 static int verify_read_ulong(void *ptr, int n)
503 {
504 int i;
505 cl_ulong *outptr = (cl_ulong *)ptr;
506
507 for (i=0; i<n; i++){
508 if( outptr[i] != TEST_PRIME_ULONG )
509 return -1;
510 }
511
512 return 0;
513 }
514
515
verify_read_short(void * ptr,int n)516 static int verify_read_short(void *ptr, int n)
517 {
518 int i;
519 short *outptr = (short *)ptr;
520
521 for (i=0; i<n; i++){
522 if( outptr[i] != (short)((1<<8)+1) )
523 return -1;
524 }
525
526 return 0;
527 }
528
529
verify_read_ushort(void * ptr,int n)530 static int verify_read_ushort(void *ptr, int n)
531 {
532 int i;
533 cl_ushort *outptr = (cl_ushort *)ptr;
534
535 for (i=0; i<n; i++){
536 if( outptr[i] != (cl_ushort)((1<<8)+1) )
537 return -1;
538 }
539
540 return 0;
541 }
542
543
verify_read_float(void * ptr,int n)544 static int verify_read_float( void *ptr, int n )
545 {
546 int i;
547 float *outptr = (float *)ptr;
548
549 for (i=0; i<n; i++){
550 if( outptr[i] != TEST_PRIME_FLOAT )
551 return -1;
552 }
553
554 return 0;
555 }
556
557
verify_read_half(void * ptr,int n)558 static int verify_read_half( void *ptr, int n )
559 {
560 int i;
561 float *outptr = (float *)ptr;
562
563 for( i = 0; i < n / 2; i++ ){
564 if( outptr[i] != TEST_PRIME_HALF )
565 return -1;
566 }
567
568 return 0;
569 }
570
571
verify_read_char(void * ptr,int n)572 static int verify_read_char(void *ptr, int n)
573 {
574 int i;
575 char *outptr = (char *)ptr;
576
577 for (i=0; i<n; i++){
578 if( outptr[i] != TEST_PRIME_CHAR )
579 return -1;
580 }
581
582 return 0;
583 }
584
585
verify_read_uchar(void * ptr,int n)586 static int verify_read_uchar( void *ptr, int n )
587 {
588 int i;
589 uchar *outptr = (uchar *)ptr;
590
591 for (i=0; i<n; i++){
592 if( outptr[i] != TEST_PRIME_CHAR )
593 return -1;
594 }
595
596 return 0;
597 }
598
599
verify_read_struct(void * ptr,int n)600 static int verify_read_struct( void *ptr, int n )
601 {
602 int i;
603 TestStruct *outptr = (TestStruct *)ptr;
604
605 for ( i = 0; i < n; i++ ){
606 if( ( outptr[i].a != TEST_PRIME_INT ) ||
607 ( outptr[i].b != TEST_PRIME_FLOAT ) )
608 return -1;
609 }
610
611 return 0;
612 }
613
614 //----- the test functions
test_stream_read(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements,size_t size,const char * type,int loops,const char * kernelCode[],const char * kernelName[],int (* fn)(void *,int))615 int test_stream_read( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements, size_t size, const char *type, int loops,
616 const char *kernelCode[], const char *kernelName[], int (*fn)(void *,int) )
617 {
618 cl_mem streams[5];
619 void *outptr[5];
620 cl_program program[5];
621 cl_kernel kernel[5];
622 cl_event readEvent;
623 cl_ulong queueStart, submitStart, readStart, readEnd;
624 size_t threads[1];
625 int err, err_count = 0;
626 int i;
627 size_t ptrSizes[5];
628
629 threads[0] = (size_t)num_elements;
630
631 ptrSizes[0] = size;
632 ptrSizes[1] = ptrSizes[0] << 1;
633 ptrSizes[2] = ptrSizes[1] << 1;
634 ptrSizes[3] = ptrSizes[2] << 1;
635 ptrSizes[4] = ptrSizes[3] << 1;
636 for( i = 0; i < loops; i++ ){
637 outptr[i] = malloc( ptrSizes[i] * num_elements );
638 if( ! outptr[i] ){
639 log_error( " unable to allocate %d bytes for outptr\n", (int)( ptrSizes[i] * num_elements ) );
640 return -1;
641 }
642 streams[i] = clCreateBuffer(context, CL_MEM_READ_WRITE,
643 ptrSizes[i] * num_elements, NULL, &err);
644 if( !streams[i] ){
645 log_error( " clCreateBuffer failed\n" );
646 free( outptr[i] );
647 return -1;
648 }
649 err = create_single_kernel_helper( context, &program[i], &kernel[i], 1, &kernelCode[i], kernelName[i] );
650 if( err ){
651 log_error( " Error creating program for %s\n", type );
652 clReleaseMemObject(streams[i]);
653 free( outptr[i] );
654 return -1;
655 }
656
657 err = clSetKernelArg( kernel[i], 0, sizeof( cl_mem ), (void *)&streams[i] );
658 if( err != CL_SUCCESS ){
659 print_error( err, "clSetKernelArg failed" );
660 clReleaseProgram( program[i] );
661 clReleaseKernel( kernel[i] );
662 clReleaseMemObject( streams[i] );
663 free( outptr[i] );
664 return -1;
665 }
666
667 err = clEnqueueNDRangeKernel( queue, kernel[i], 1, NULL, threads, NULL, 0, NULL, NULL );
668
669 if( err != CL_SUCCESS ){
670 print_error( err, "clEnqueueNDRangeKernel failed" );
671 clReleaseKernel( kernel[i] );
672 clReleaseProgram( program[i] );
673 clReleaseMemObject( streams[i] );
674 free( outptr[i] );
675 return -1;
676 }
677
678 err = clEnqueueReadBuffer( queue, streams[i], false, 0, ptrSizes[i]*num_elements, outptr[i], 0, NULL, &readEvent );
679 if( err != CL_SUCCESS ){
680 print_error( err, "clEnqueueReadBuffer failed" );
681 clReleaseKernel( kernel[i] );
682 clReleaseProgram( program[i] );
683 clReleaseMemObject( streams[i] );
684 free( outptr[i] );
685 return -1;
686 }
687 err = clWaitForEvents( 1, &readEvent );
688 if( err != CL_SUCCESS )
689 {
690 print_error( err, "Unable to wait for event completion" );
691 clReleaseKernel( kernel[i] );
692 clReleaseProgram( program[i] );
693 clReleaseMemObject( streams[i] );
694 free( outptr[i] );
695 return -1;
696 }
697 err = clGetEventProfilingInfo( readEvent, CL_PROFILING_COMMAND_QUEUED, sizeof( cl_ulong ), &queueStart, NULL );
698 if( err != CL_SUCCESS ){
699 print_error( err, "clGetEventProfilingInfo failed" );
700 clReleaseKernel( kernel[i] );
701 clReleaseProgram( program[i] );
702 clReleaseMemObject( streams[i] );
703 free( outptr[i] );
704 return -1;
705 }
706
707 err = clGetEventProfilingInfo( readEvent, CL_PROFILING_COMMAND_SUBMIT, sizeof( cl_ulong ), &submitStart, NULL );
708 if( err != CL_SUCCESS ){
709 print_error( err, "clGetEventProfilingInfo failed" );
710 clReleaseKernel( kernel[i] );
711 clReleaseProgram( program[i] );
712 clReleaseMemObject( streams[i] );
713 free( outptr[i] );
714 return -1;
715 }
716
717 err = clGetEventProfilingInfo( readEvent, CL_PROFILING_COMMAND_START, sizeof( cl_ulong ), &readStart, NULL );
718 if( err != CL_SUCCESS ){
719 print_error( err, "clGetEventProfilingInfo failed" );
720 clReleaseKernel( kernel[i] );
721 clReleaseProgram( program[i] );
722 clReleaseMemObject( streams[i] );
723 free( outptr[i] );
724 return -1;
725 }
726
727 err = clGetEventProfilingInfo( readEvent, CL_PROFILING_COMMAND_END, sizeof( cl_ulong ), &readEnd, NULL );
728 if( err != CL_SUCCESS ){
729 print_error( err, "clGetEventProfilingInfo failed" );
730 clReleaseKernel( kernel[i] );
731 clReleaseProgram( program[i] );
732 clReleaseMemObject( streams[i] );
733 free( outptr[i] );
734 return -1;
735 }
736
737 if (fn(outptr[i], num_elements*(1<<i))){
738 log_error( " %s%d data failed to verify\n", type, 1<<i );
739 err_count++;
740 }
741 else{
742 log_info( " %s%d data verified\n", type, 1<<i );
743 }
744
745 if (check_times(queueStart, submitStart, readStart, readEnd, device))
746 err_count++;
747
748 // cleanup
749 clReleaseEvent(readEvent);
750 clReleaseKernel( kernel[i] );
751 clReleaseProgram( program[i] );
752 clReleaseMemObject( streams[i] );
753 free( outptr[i] );
754 }
755
756 return err_count;
757
758 } // end test_stream_read()
759
760
test_read_array_int(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)761 int test_read_array_int( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
762 {
763 int (*foo)(void *,int);
764 foo = verify_read_int;
765
766 return test_stream_read( device, context, queue, num_elements, sizeof( cl_int ), "int", 5,
767 stream_read_int_kernel_code, int_kernel_name, foo );
768 }
769
770
test_read_array_uint(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)771 int test_read_array_uint( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
772 {
773 int (*foo)(void *,int);
774 foo = verify_read_uint;
775
776 return test_stream_read( device, context, queue, num_elements, sizeof( cl_uint ), "uint", 5,
777 stream_read_uint_kernel_code, uint_kernel_name, foo );
778 }
779
780
test_read_array_long(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)781 int test_read_array_long( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
782 {
783 int (*foo)(void *,int);
784 foo = verify_read_long;
785
786 if (!gHasLong)
787 {
788 log_info("read_long_array: Long types unsupported, skipping.");
789 return CL_SUCCESS;
790 }
791
792 return test_stream_read( device, context, queue, num_elements, sizeof( cl_long ), "long", 5,
793 stream_read_long_kernel_code, long_kernel_name, foo );
794 }
795
796
test_read_array_ulong(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)797 int test_read_array_ulong( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
798 {
799 int (*foo)(void *,int);
800 foo = verify_read_ulong;
801
802 if (!gHasLong)
803 {
804 log_info("read_long_array: Long types unsupported, skipping.");
805 return CL_SUCCESS;
806 }
807
808 return test_stream_read( device, context, queue, num_elements, sizeof( cl_ulong ), "ulong", 5,
809 stream_read_ulong_kernel_code, ulong_kernel_name, foo );
810 }
811
812
test_read_array_short(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)813 int test_read_array_short( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
814 {
815 int (*foo)(void *,int);
816 foo = verify_read_short;
817
818 return test_stream_read( device, context, queue, num_elements, sizeof( cl_short ), "short", 5,
819 stream_read_short_kernel_code, short_kernel_name, foo );
820 }
821
822
test_read_array_ushort(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)823 int test_read_array_ushort( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
824 {
825 int (*foo)(void *,int);
826 foo = verify_read_ushort;
827
828 return test_stream_read( device, context, queue, num_elements, sizeof( cl_ushort ), "ushort", 5,
829 stream_read_ushort_kernel_code, ushort_kernel_name, foo );
830 }
831
832
test_read_array_float(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)833 int test_read_array_float( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
834 {
835 int (*foo)(void *,int);
836 foo = verify_read_float;
837
838 return test_stream_read( device, context, queue, num_elements, sizeof( cl_float ), "float", 5,
839 stream_read_float_kernel_code, float_kernel_name, foo );
840 }
841
842
test_read_array_half(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)843 int test_read_array_half( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
844 {
845 int (*foo)(void *,int);
846 foo = verify_read_half;
847
848 return test_stream_read( device, context, queue, num_elements, sizeof( cl_half ), "half", 5,
849 stream_read_half_kernel_code, half_kernel_name, foo );
850 }
851
852
test_read_array_char(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)853 int test_read_array_char( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
854 {
855 int (*foo)(void *,int);
856 foo = verify_read_char;
857
858 return test_stream_read( device, context, queue, num_elements, sizeof( cl_char ), "char", 5,
859 stream_read_char_kernel_code, char_kernel_name, foo );
860 }
861
862
test_read_array_uchar(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)863 int test_read_array_uchar( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
864 {
865 int (*foo)(void *,int);
866 foo = verify_read_uchar;
867
868 return test_stream_read( device, context, queue, num_elements, sizeof( cl_uchar ), "uchar", 5,
869 stream_read_uchar_kernel_code, uchar_kernel_name, foo );
870 }
871
872
test_read_array_struct(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)873 int test_read_array_struct( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
874 {
875 int (*foo)(void *,int);
876 foo = verify_read_struct;
877
878 return test_stream_read( device, context, queue, num_elements, sizeof( TestStruct ), "struct", 1,
879 stream_read_struct_kernel_code, struct_kernel_name, foo );
880 }
881