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 #define TEST_PRIME_INT ((1<<16)+1)
29 #define TEST_PRIME_UINT ((1U<<16)+1U)
30 #define TEST_PRIME_LONG ((1LL<<32)+1LL)
31 #define TEST_PRIME_ULONG ((1ULL<<32)+1ULL)
32 #define TEST_PRIME_SHORT ((1S<<8)+1S)
33 #define TEST_PRIME_FLOAT (float)3.40282346638528860e+38
34 #define TEST_PRIME_HALF 119.f
35 #define TEST_BOOL true
36 #define TEST_PRIME_CHAR 0x77
37
38
39 #ifndef TestStruct
40 typedef struct{
41 int a;
42 float b;
43 } TestStruct;
44 #endif
45
46
47 //--- the code for the kernel executables
48 static const char *buffer_read_int_kernel_code[] = {
49 "__kernel void test_buffer_read_int(__global int *dst)\n"
50 "{\n"
51 " int tid = get_global_id(0);\n"
52 "\n"
53 " dst[tid] = ((1<<16)+1);\n"
54 "}\n",
55
56 "__kernel void test_buffer_read_int2(__global int2 *dst)\n"
57 "{\n"
58 " int tid = get_global_id(0);\n"
59 "\n"
60 " dst[tid] = ((1<<16)+1);\n"
61 "}\n",
62
63 "__kernel void test_buffer_read_int4(__global int4 *dst)\n"
64 "{\n"
65 " int tid = get_global_id(0);\n"
66 "\n"
67 " dst[tid] = ((1<<16)+1);\n"
68 "}\n",
69
70 "__kernel void test_buffer_read_int8(__global int8 *dst)\n"
71 "{\n"
72 " int tid = get_global_id(0);\n"
73 "\n"
74 " dst[tid] = ((1<<16)+1);\n"
75 "}\n",
76
77 "__kernel void test_buffer_read_int16(__global int16 *dst)\n"
78 "{\n"
79 " int tid = get_global_id(0);\n"
80 "\n"
81 " dst[tid] = ((1<<16)+1);\n"
82 "}\n" };
83
84 static const char *int_kernel_name[] = { "test_buffer_read_int", "test_buffer_read_int2", "test_buffer_read_int4", "test_buffer_read_int8", "test_buffer_read_int16" };
85
86 static const char *buffer_read_uint_kernel_code[] = {
87 "__kernel void test_buffer_read_uint(__global uint *dst)\n"
88 "{\n"
89 " int tid = get_global_id(0);\n"
90 "\n"
91 " dst[tid] = ((1U<<16)+1U);\n"
92 "}\n",
93
94 "__kernel void test_buffer_read_uint2(__global uint2 *dst)\n"
95 "{\n"
96 " int tid = get_global_id(0);\n"
97 "\n"
98 " dst[tid] = ((1U<<16)+1U);\n"
99 "}\n",
100
101 "__kernel void test_buffer_read_uint4(__global uint4 *dst)\n"
102 "{\n"
103 " int tid = get_global_id(0);\n"
104 "\n"
105 " dst[tid] = ((1U<<16)+1U);\n"
106 "}\n",
107
108 "__kernel void test_buffer_read_uint8(__global uint8 *dst)\n"
109 "{\n"
110 " int tid = get_global_id(0);\n"
111 "\n"
112 " dst[tid] = ((1U<<16)+1U);\n"
113 "}\n",
114
115 "__kernel void test_buffer_read_uint16(__global uint16 *dst)\n"
116 "{\n"
117 " int tid = get_global_id(0);\n"
118 "\n"
119 " dst[tid] = ((1U<<16)+1U);\n"
120 "}\n" };
121
122 static const char *uint_kernel_name[] = { "test_buffer_read_uint", "test_buffer_read_uint2", "test_buffer_read_uint4", "test_buffer_read_uint8", "test_buffer_read_uint16" };
123
124 static const char *buffer_read_long_kernel_code[] = {
125 "__kernel void test_buffer_read_long(__global long *dst)\n"
126 "{\n"
127 " int tid = get_global_id(0);\n"
128 "\n"
129 " dst[tid] = ((1L<<32)+1L);\n"
130 "}\n",
131
132 "__kernel void test_buffer_read_long2(__global long2 *dst)\n"
133 "{\n"
134 " int tid = get_global_id(0);\n"
135 "\n"
136 " dst[tid] = ((1L<<32)+1L);\n"
137 "}\n",
138
139 "__kernel void test_buffer_read_long4(__global long4 *dst)\n"
140 "{\n"
141 " int tid = get_global_id(0);\n"
142 "\n"
143 " dst[tid] = ((1L<<32)+1L);\n"
144 "}\n",
145
146 "__kernel void test_buffer_read_long8(__global long8 *dst)\n"
147 "{\n"
148 " int tid = get_global_id(0);\n"
149 "\n"
150 " dst[tid] = ((1L<<32)+1L);\n"
151 "}\n",
152
153 "__kernel void test_buffer_read_long16(__global long16 *dst)\n"
154 "{\n"
155 " int tid = get_global_id(0);\n"
156 "\n"
157 " dst[tid] = ((1L<<32)+1L);\n"
158 "}\n" };
159
160 static const char *long_kernel_name[] = { "test_buffer_read_long", "test_buffer_read_long2", "test_buffer_read_long4", "test_buffer_read_long8", "test_buffer_read_long16" };
161
162 static const char *buffer_read_ulong_kernel_code[] = {
163 "__kernel void test_buffer_read_ulong(__global ulong *dst)\n"
164 "{\n"
165 " int tid = get_global_id(0);\n"
166 "\n"
167 " dst[tid] = ((1UL<<32)+1UL);\n"
168 "}\n",
169
170 "__kernel void test_buffer_read_ulong2(__global ulong2 *dst)\n"
171 "{\n"
172 " int tid = get_global_id(0);\n"
173 "\n"
174 " dst[tid] = ((1UL<<32)+1UL);\n"
175 "}\n",
176
177 "__kernel void test_buffer_read_ulong4(__global ulong4 *dst)\n"
178 "{\n"
179 " int tid = get_global_id(0);\n"
180 "\n"
181 " dst[tid] = ((1UL<<32)+1UL);\n"
182 "}\n",
183
184 "__kernel void test_buffer_read_ulong8(__global ulong8 *dst)\n"
185 "{\n"
186 " int tid = get_global_id(0);\n"
187 "\n"
188 " dst[tid] = ((1UL<<32)+1UL);\n"
189 "}\n",
190
191 "__kernel void test_buffer_read_ulong16(__global ulong16 *dst)\n"
192 "{\n"
193 " int tid = get_global_id(0);\n"
194 "\n"
195 " dst[tid] = ((1UL<<32)+1UL);\n"
196 "}\n" };
197
198 static const char *ulong_kernel_name[] = { "test_buffer_read_ulong", "test_buffer_read_ulong2", "test_buffer_read_ulong4", "test_buffer_read_ulong8", "test_buffer_read_ulong16" };
199
200 static const char *buffer_read_short_kernel_code[] = {
201 "__kernel void test_buffer_read_short(__global short *dst)\n"
202 "{\n"
203 " int tid = get_global_id(0);\n"
204 "\n"
205 " dst[tid] = (short)((1<<8)+1);\n"
206 "}\n",
207
208 "__kernel void test_buffer_read_short2(__global short2 *dst)\n"
209 "{\n"
210 " int tid = get_global_id(0);\n"
211 "\n"
212 " dst[tid] = (short)((1<<8)+1);\n"
213 "}\n",
214
215 "__kernel void test_buffer_read_short4(__global short4 *dst)\n"
216 "{\n"
217 " int tid = get_global_id(0);\n"
218 "\n"
219 " dst[tid] = (short)((1<<8)+1);\n"
220 "}\n",
221
222 "__kernel void test_buffer_read_short8(__global short8 *dst)\n"
223 "{\n"
224 " int tid = get_global_id(0);\n"
225 "\n"
226 " dst[tid] = (short)((1<<8)+1);\n"
227 "}\n",
228
229 "__kernel void test_buffer_read_short16(__global short16 *dst)\n"
230 "{\n"
231 " int tid = get_global_id(0);\n"
232 "\n"
233 " dst[tid] = (short)((1<<8)+1);\n"
234 "}\n" };
235
236 static const char *short_kernel_name[] = { "test_buffer_read_short", "test_buffer_read_short2", "test_buffer_read_short4", "test_buffer_read_short8", "test_buffer_read_short16" };
237
238
239 static const char *buffer_read_ushort_kernel_code[] = {
240 "__kernel void test_buffer_read_ushort(__global ushort *dst)\n"
241 "{\n"
242 " int tid = get_global_id(0);\n"
243 "\n"
244 " dst[tid] = (ushort)((1<<8)+1);\n"
245 "}\n",
246
247 "__kernel void test_buffer_read_ushort2(__global ushort2 *dst)\n"
248 "{\n"
249 " int tid = get_global_id(0);\n"
250 "\n"
251 " dst[tid] = (ushort)((1<<8)+1);\n"
252 "}\n",
253
254 "__kernel void test_buffer_read_ushort4(__global ushort4 *dst)\n"
255 "{\n"
256 " int tid = get_global_id(0);\n"
257 "\n"
258 " dst[tid] = (ushort)((1<<8)+1);\n"
259 "}\n",
260
261 "__kernel void test_buffer_read_ushort8(__global ushort8 *dst)\n"
262 "{\n"
263 " int tid = get_global_id(0);\n"
264 "\n"
265 " dst[tid] = (ushort)((1<<8)+1);\n"
266 "}\n",
267
268 "__kernel void test_buffer_read_ushort16(__global ushort16 *dst)\n"
269 "{\n"
270 " int tid = get_global_id(0);\n"
271 "\n"
272 " dst[tid] = (ushort)((1<<8)+1);\n"
273 "}\n" };
274
275 static const char *ushort_kernel_name[] = { "test_buffer_read_ushort", "test_buffer_read_ushort2", "test_buffer_read_ushort4", "test_buffer_read_ushort8", "test_buffer_read_ushort16" };
276
277
278 static const char *buffer_read_float_kernel_code[] = {
279 "__kernel void test_buffer_read_float(__global float *dst)\n"
280 "{\n"
281 " int tid = get_global_id(0);\n"
282 "\n"
283 " dst[tid] = (float)3.40282346638528860e+38;\n"
284 "}\n",
285
286 "__kernel void test_buffer_read_float2(__global float2 *dst)\n"
287 "{\n"
288 " int tid = get_global_id(0);\n"
289 "\n"
290 " dst[tid] = (float)3.40282346638528860e+38;\n"
291 "}\n",
292
293 "__kernel void test_buffer_read_float4(__global float4 *dst)\n"
294 "{\n"
295 " int tid = get_global_id(0);\n"
296 "\n"
297 " dst[tid] = (float)3.40282346638528860e+38;\n"
298 "}\n",
299
300 "__kernel void test_buffer_read_float8(__global float8 *dst)\n"
301 "{\n"
302 " int tid = get_global_id(0);\n"
303 "\n"
304 " dst[tid] = (float)3.40282346638528860e+38;\n"
305 "}\n",
306
307 "__kernel void test_buffer_read_float16(__global float16 *dst)\n"
308 "{\n"
309 " int tid = get_global_id(0);\n"
310 "\n"
311 " dst[tid] = (float)3.40282346638528860e+38;\n"
312 "}\n" };
313
314 static const char *float_kernel_name[] = { "test_buffer_read_float", "test_buffer_read_float2", "test_buffer_read_float4", "test_buffer_read_float8", "test_buffer_read_float16" };
315
316
317 static const char *buffer_read_char_kernel_code[] = {
318 "__kernel void test_buffer_read_char(__global char *dst)\n"
319 "{\n"
320 " int tid = get_global_id(0);\n"
321 "\n"
322 " dst[tid] = (char)'w';\n"
323 "}\n",
324
325 "__kernel void test_buffer_read_char2(__global char2 *dst)\n"
326 "{\n"
327 " int tid = get_global_id(0);\n"
328 "\n"
329 " dst[tid] = (char)'w';\n"
330 "}\n",
331
332 "__kernel void test_buffer_read_char4(__global char4 *dst)\n"
333 "{\n"
334 " int tid = get_global_id(0);\n"
335 "\n"
336 " dst[tid] = (char)'w';\n"
337 "}\n",
338
339 "__kernel void test_buffer_read_char8(__global char8 *dst)\n"
340 "{\n"
341 " int tid = get_global_id(0);\n"
342 "\n"
343 " dst[tid] = (char)'w';\n"
344 "}\n",
345
346 "__kernel void test_buffer_read_char16(__global char16 *dst)\n"
347 "{\n"
348 " int tid = get_global_id(0);\n"
349 "\n"
350 " dst[tid] = (char)'w';\n"
351 "}\n" };
352
353 static const char *char_kernel_name[] = { "test_buffer_read_char", "test_buffer_read_char2", "test_buffer_read_char4", "test_buffer_read_char8", "test_buffer_read_char16" };
354
355
356 static const char *buffer_read_uchar_kernel_code[] = {
357 "__kernel void test_buffer_read_uchar(__global uchar *dst)\n"
358 "{\n"
359 " int tid = get_global_id(0);\n"
360 "\n"
361 " dst[tid] = 'w';\n"
362 "}\n",
363
364 "__kernel void test_buffer_read_uchar2(__global uchar2 *dst)\n"
365 "{\n"
366 " int tid = get_global_id(0);\n"
367 "\n"
368 " dst[tid] = (uchar)'w';\n"
369 "}\n",
370
371 "__kernel void test_buffer_read_uchar4(__global uchar4 *dst)\n"
372 "{\n"
373 " int tid = get_global_id(0);\n"
374 "\n"
375 " dst[tid] = (uchar)'w';\n"
376 "}\n",
377
378 "__kernel void test_buffer_read_uchar8(__global uchar8 *dst)\n"
379 "{\n"
380 " int tid = get_global_id(0);\n"
381 "\n"
382 " dst[tid] = (uchar)'w';\n"
383 "}\n",
384
385 "__kernel void test_buffer_read_uchar16(__global uchar16 *dst)\n"
386 "{\n"
387 " int tid = get_global_id(0);\n"
388 "\n"
389 " dst[tid] = (uchar)'w';\n"
390 "}\n" };
391
392 static const char *uchar_kernel_name[] = { "test_buffer_read_uchar", "test_buffer_read_uchar2", "test_buffer_read_uchar4", "test_buffer_read_uchar8", "test_buffer_read_uchar16" };
393
394
395 static const char *buffer_read_struct_kernel_code[] = {
396 "typedef struct{\n"
397 "int a;\n"
398 "float b;\n"
399 "} TestStruct;\n"
400 "__kernel void test_buffer_read_struct(__global TestStruct *dst)\n"
401 "{\n"
402 " int tid = get_global_id(0);\n"
403 "\n"
404 " dst[tid].a = ((1<<16)+1);\n"
405 " dst[tid].b = (float)3.40282346638528860e+38;\n"
406 "}\n" };
407
408 static const char *struct_kernel_name[] = { "test_buffer_read_struct" };
409
410
411 //--- the verify functions
verify_read_int(void * ptr,int n)412 static int verify_read_int(void *ptr, int n)
413 {
414 int i;
415 int *outptr = (int *)ptr;
416
417 for (i=0; i<n; i++){
418 if ( outptr[i] != TEST_PRIME_INT )
419 return -1;
420 }
421
422 return 0;
423 }
424
425
verify_read_uint(void * ptr,int n)426 static int verify_read_uint(void *ptr, int n)
427 {
428 int i;
429 cl_uint *outptr = (cl_uint *)ptr;
430
431 for (i=0; i<n; i++){
432 if ( outptr[i] != TEST_PRIME_UINT )
433 return -1;
434 }
435
436 return 0;
437 }
438
439
verify_read_long(void * ptr,int n)440 static int verify_read_long(void *ptr, int n)
441 {
442 int i;
443 cl_long *outptr = (cl_long *)ptr;
444
445 for (i=0; i<n; i++){
446 if ( outptr[i] != TEST_PRIME_LONG )
447 return -1;
448 }
449
450 return 0;
451 }
452
453
verify_read_ulong(void * ptr,int n)454 static int verify_read_ulong(void *ptr, int n)
455 {
456 int i;
457 cl_ulong *outptr = (cl_ulong *)ptr;
458
459 for (i=0; i<n; i++){
460 if ( outptr[i] != TEST_PRIME_ULONG )
461 return -1;
462 }
463
464 return 0;
465 }
466
467
verify_read_short(void * ptr,int n)468 static int verify_read_short(void *ptr, int n)
469 {
470 int i;
471 short *outptr = (short *)ptr;
472
473 for (i=0; i<n; i++){
474 if ( outptr[i] != (short)((1<<8)+1) )
475 return -1;
476 }
477
478 return 0;
479 }
480
481
verify_read_ushort(void * ptr,int n)482 static int verify_read_ushort(void *ptr, int n)
483 {
484 int i;
485 cl_ushort *outptr = (cl_ushort *)ptr;
486
487 for (i=0; i<n; i++){
488 if ( outptr[i] != (cl_ushort)((1<<8)+1) )
489 return -1;
490 }
491
492 return 0;
493 }
494
495
verify_read_float(void * ptr,int n)496 static int verify_read_float( void *ptr, int n )
497 {
498 int i;
499 float *outptr = (float *)ptr;
500
501 for (i=0; i<n; i++){
502 if ( outptr[i] != TEST_PRIME_FLOAT )
503 return -1;
504 }
505
506 return 0;
507 }
508
509
verify_read_char(void * ptr,int n)510 static int verify_read_char(void *ptr, int n)
511 {
512 int i;
513 char *outptr = (char *)ptr;
514
515 for (i=0; i<n; i++){
516 if ( outptr[i] != TEST_PRIME_CHAR )
517 return -1;
518 }
519
520 return 0;
521 }
522
523
verify_read_uchar(void * ptr,int n)524 static int verify_read_uchar( void *ptr, int n )
525 {
526 int i;
527 cl_uchar *outptr = (cl_uchar *)ptr;
528
529 for ( i = 0; i < n; i++ ){
530 if ( outptr[i] != TEST_PRIME_CHAR )
531 return -1;
532 }
533
534 return 0;
535 }
536
537
verify_read_struct(void * ptr,int n)538 static int verify_read_struct( void *ptr, int n )
539 {
540 int i;
541 TestStruct *outptr = (TestStruct *)ptr;
542
543 for ( i = 0; i < n; i++ ){
544 if ( ( outptr[i].a != TEST_PRIME_INT ) ||
545 ( outptr[i].b != TEST_PRIME_FLOAT ) )
546 return -1;
547 }
548
549 return 0;
550 }
551
552
553 //----- the test functions
test_buffer_map_read(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements,size_t size,char * type,int loops,const char * kernelCode[],const char * kernelName[],int (* fn)(void *,int))554 static int test_buffer_map_read( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, size_t size, char *type, int loops,
555 const char *kernelCode[], const char *kernelName[], int (*fn)(void *,int) )
556 {
557 void *outptr[5];
558 clProgramWrapper program[5];
559 clKernelWrapper kernel[5];
560 size_t threads[3], localThreads[3];
561 cl_int err;
562 int i;
563 size_t ptrSizes[5];
564 int src_flag_id;
565 int total_errors = 0;
566 void *mappedPtr;
567
568 size_t min_alignment = get_min_alignment(context);
569
570 threads[0] = (cl_uint)num_elements;
571
572 ptrSizes[0] = size;
573 ptrSizes[1] = ptrSizes[0] << 1;
574 ptrSizes[2] = ptrSizes[1] << 1;
575 ptrSizes[3] = ptrSizes[2] << 1;
576 ptrSizes[4] = ptrSizes[3] << 1;
577
578 //embedded devices don't support long/ulong so skip over
579 if (! gHasLong && strstr(type,"long"))
580 return 0;
581
582 for (i = 0; i < loops; i++)
583 {
584
585 err = create_single_kernel_helper(context, &program[i], &kernel[i], 1,
586 &kernelCode[i], kernelName[i]);
587 if (err)
588 {
589 log_error(" Error creating program for %s\n", type);
590 return -1;
591 }
592
593 for (src_flag_id = 0; src_flag_id < NUM_FLAGS; src_flag_id++)
594 {
595 clMemWrapper buffer;
596 outptr[i] = align_malloc( ptrSizes[i] * num_elements, min_alignment);
597 if ( ! outptr[i] ){
598 log_error( " unable to allocate %d bytes of memory\n", (int)ptrSizes[i] * num_elements );
599 return -1;
600 }
601
602 if ((flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR) || (flag_set[src_flag_id] & CL_MEM_COPY_HOST_PTR))
603 buffer =
604 clCreateBuffer(context, flag_set[src_flag_id],
605 ptrSizes[i] * num_elements, outptr[i], &err);
606 else
607 buffer = clCreateBuffer(context, flag_set[src_flag_id],
608 ptrSizes[i] * num_elements, NULL, &err);
609
610 if (!buffer || err)
611 {
612 print_error(err, "clCreateBuffer failed\n" );
613 align_free( outptr[i] );
614 return -1;
615 }
616
617 err = clSetKernelArg(kernel[i], 0, sizeof(cl_mem), (void *)&buffer);
618
619 if ( err != CL_SUCCESS ){
620 print_error( err, "clSetKernelArg failed\n" );
621 align_free( outptr[i] );
622 return -1;
623 }
624
625 threads[0] = (cl_uint)num_elements;
626
627 err = get_max_common_work_group_size( context, kernel[i], threads[0], &localThreads[0] );
628 test_error( err, "Unable to get work group size to use" );
629
630 err = clEnqueueNDRangeKernel( queue, kernel[i], 1, NULL, threads, localThreads, 0, NULL, NULL );
631 if ( err != CL_SUCCESS ){
632 print_error( err, "clEnqueueNDRangeKernel failed\n" );
633 align_free( outptr[i] );
634 return -1;
635 }
636
637 mappedPtr = clEnqueueMapBuffer(queue, buffer, CL_TRUE, CL_MAP_READ,
638 0, ptrSizes[i] * num_elements, 0,
639 NULL, NULL, &err);
640 if (err != CL_SUCCESS)
641 {
642 print_error( err, "clEnqueueMapBuffer failed" );
643 align_free( outptr[i] );
644 return -1;
645 }
646
647 if (fn(mappedPtr, num_elements*(1<<i))){
648 log_error(" %s%d test failed. cl_mem_flags src: %s\n", type,
649 1 << i, flag_set_names[src_flag_id]);
650 total_errors++;
651 }
652 else{
653 log_info(" %s%d test passed. cl_mem_flags src: %s\n", type,
654 1 << i, flag_set_names[src_flag_id]);
655 }
656
657 err = clEnqueueUnmapMemObject(queue, buffer, mappedPtr, 0, NULL,
658 NULL);
659 test_error(err, "clEnqueueUnmapMemObject failed");
660
661 // If we are using the outptr[i] as backing via USE_HOST_PTR we need to make sure we are done before freeing.
662 if ((flag_set[src_flag_id] & CL_MEM_USE_HOST_PTR)) {
663 err = clFinish(queue);
664 test_error(err, "clFinish failed");
665 }
666 align_free( outptr[i] );
667 }
668 } // cl_mem_flags
669
670 return total_errors;
671
672 } // end test_buffer_map_read()
673
674
675 #define DECLARE_LOCK_TEST(type, realType) \
676 int test_buffer_map_read_##type( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements ) \
677 { \
678 return test_buffer_map_read( deviceID, context, queue, num_elements, sizeof( realType ), (char*)#type, 5, \
679 buffer_read_##type##_kernel_code, type##_kernel_name, verify_read_##type ); \
680 }
681
DECLARE_LOCK_TEST(int,cl_int)682 DECLARE_LOCK_TEST(int, cl_int)
683 DECLARE_LOCK_TEST(uint, cl_uint)
684 DECLARE_LOCK_TEST(long, cl_long)
685 DECLARE_LOCK_TEST(ulong, cl_ulong)
686 DECLARE_LOCK_TEST(short, cl_short)
687 DECLARE_LOCK_TEST(ushort, cl_ushort)
688 DECLARE_LOCK_TEST(char, cl_char)
689 DECLARE_LOCK_TEST(uchar, cl_uchar)
690 DECLARE_LOCK_TEST(float, cl_float)
691
692 int test_buffer_map_read_struct( cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements )
693 {
694 int (*foo)(void *,int);
695 foo = verify_read_struct;
696
697 return test_buffer_map_read( deviceID, context, queue, num_elements, sizeof( TestStruct ), (char*)"struct", 1,
698 buffer_read_struct_kernel_code, struct_kernel_name, foo );
699
700 } // end test_buffer_map_struct_read()
701
702