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 <sys/types.h>
21 #include <sys/stat.h>
22
23 #include "procs.h"
24 #include "harness/testHarness.h"
25 #include "harness/errorHelpers.h"
26 #include "harness/conversions.h"
27
28 //--- the code for the kernel executables
29 static const char *write_kernel_code =
30 "\n"
31 "__kernel void test_write(__global unsigned char *src, write_only image2d_t dstimg)\n"
32 "{\n"
33 " int tid_x = get_global_id(0);\n"
34 " int tid_y = get_global_id(1);\n"
35 " int indx = tid_y * get_image_width(dstimg) + tid_x;\n"
36 " float4 color;\n"
37 "\n"
38 " indx *= 4;\n"
39 " color = (float4)((float)src[indx+0], (float)src[indx+1], (float)src[indx+2], (float)src[indx+3]);\n"
40 " color /= (float4)(255.0f, 255.0f, 255.0f, 255.0f);\n"
41 " write_imagef(dstimg, (int2)(tid_x, tid_y), color);\n"
42 "\n"
43 "}\n";
44
45
46 //--- the verify functions
verify_subimage(unsigned char * src,unsigned char * dst,size_t srcx,size_t srcy,size_t dstx,size_t dsty,size_t subw,size_t subh,size_t pitch,size_t element_pitch)47 static int verify_subimage( unsigned char *src, unsigned char *dst, size_t srcx, size_t srcy,
48 size_t dstx, size_t dsty, size_t subw, size_t subh, size_t pitch, size_t element_pitch )
49 {
50 size_t i, j, k;
51 size_t srcj, dstj;
52 size_t srcLoc, dstLoc;
53
54 for( j = 0; j < subh; j++ ){
55 srcj = ( j + srcy ) * pitch * element_pitch;
56 dstj = ( j + dsty ) * pitch * element_pitch;
57 for( i = 0; i < subw; i++ ){
58 srcLoc = srcj + ( i + srcx ) * element_pitch;
59 dstLoc = dstj + ( i + dstx ) * element_pitch;
60 for( k = 0; k < element_pitch; k++ ){ // test each channel
61 if( src[srcLoc+k] != dst[dstLoc+k] ){
62 return -1;
63 }
64 }
65 }
66 }
67
68 return 0;
69 }
70
71
verify_copy_array(int * inptr,int * outptr,int n)72 static int verify_copy_array( int *inptr, int *outptr, int n )
73 {
74 int i;
75
76 for( i = 0; i < n; i++ ) {
77 if( outptr[i] != inptr[i] )
78 return -1;
79 }
80
81 return 0;
82 }
83
84
85 //----- helper functions
generate_image(int n,MTdata d)86 static cl_uchar *generate_image( int n, MTdata d )
87 {
88 cl_uchar *ptr = (cl_uchar *)malloc( n );
89 int i;
90
91 for( i = 0; i < n; i++ )
92 ptr[i] = (cl_uchar)genrand_int32(d);
93
94 return ptr;
95 }
96
97
copy_size(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements,MTdata d)98 static int copy_size( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements, MTdata d )
99 {
100 cl_mem streams[2];
101 cl_event copyEvent;
102 cl_ulong queueStart, submitStart, writeStart, writeEnd;
103 cl_int *int_input_ptr, *int_output_ptr;
104 int err = 0;
105 int i;
106
107 int_input_ptr = (cl_int*)malloc(sizeof(cl_int) * num_elements);
108 int_output_ptr = (cl_int*)malloc(sizeof(cl_int) * num_elements);
109
110 streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
111 sizeof(cl_int) * num_elements, NULL, &err);
112 if( !streams[0] ){
113 log_error("clCreateBuffer failed\n");
114 return -1;
115 }
116 streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
117 sizeof(cl_int) * num_elements, NULL, &err);
118 if( !streams[1] ){
119 log_error("clCreateBuffer failed\n");
120 return -1;
121 }
122
123 for (i=0; i<num_elements; i++){
124 int_input_ptr[i] = (int)genrand_int32(d);
125 int_output_ptr[i] = (int)genrand_int32(d) >> 30; // seed with incorrect data
126 }
127
128 err = clEnqueueWriteBuffer( queue, streams[0], true, 0, sizeof(cl_int)*num_elements, (void *)int_input_ptr, 0, NULL, NULL );
129 if( err != CL_SUCCESS ){
130 print_error( err, "clWriteArray failed" );
131 clReleaseMemObject( streams[0] );
132 clReleaseMemObject( streams[1] );
133 free( (void *)int_output_ptr );
134 free( (void *)int_input_ptr );
135 return -1;
136 }
137
138 err = clEnqueueCopyBuffer( queue, streams[0], streams[1], 0, 0, sizeof(cl_int)*num_elements, 0, NULL, ©Event );
139 if( err != CL_SUCCESS ){
140 print_error( err, "clCopyArray failed" );
141 clReleaseMemObject( streams[0] );
142 clReleaseMemObject( streams[1] );
143 free( (void *)int_output_ptr );
144 free( (void *)int_input_ptr );
145 return -1;
146 }
147
148 // This synchronization point is needed in order to assume the data is valid.
149 // Getting profiling information is not a synchronization point.
150 err = clWaitForEvents( 1, ©Event );
151 if( err != CL_SUCCESS )
152 {
153 clReleaseEvent(copyEvent);
154 clReleaseMemObject( streams[0] );
155 clReleaseMemObject( streams[1] );
156 free( (void *)int_output_ptr );
157 free( (void *)int_input_ptr );
158 return -1;
159 }
160
161 // test profiling
162 while( ( err = clGetEventProfilingInfo( copyEvent, CL_PROFILING_COMMAND_QUEUED, sizeof( cl_ulong ), &queueStart, NULL ) ) ==
163 CL_PROFILING_INFO_NOT_AVAILABLE );
164 if( err != CL_SUCCESS ){
165 print_error( err, "clGetEventProfilingInfo failed" );
166 clReleaseEvent(copyEvent);
167 clReleaseMemObject( streams[0] );
168 clReleaseMemObject( streams[1] );
169 free( (void *)int_output_ptr );
170 free( (void *)int_input_ptr );
171 return -1;
172 }
173
174 while( ( err = clGetEventProfilingInfo( copyEvent, CL_PROFILING_COMMAND_SUBMIT, sizeof( cl_ulong ), &submitStart, NULL ) ) ==
175 CL_PROFILING_INFO_NOT_AVAILABLE );
176 if( err != CL_SUCCESS ){
177 print_error( err, "clGetEventProfilingInfo failed" );
178 clReleaseEvent(copyEvent);
179 clReleaseMemObject( streams[0] );
180 clReleaseMemObject( streams[1] );
181 free( (void *)int_output_ptr );
182 free( (void *)int_input_ptr );
183 return -1;
184 }
185
186 err = clGetEventProfilingInfo( copyEvent, CL_PROFILING_COMMAND_START, sizeof( cl_ulong ), &writeStart, NULL );
187 if( err != CL_SUCCESS ){
188 print_error( err, "clGetEventProfilingInfo failed" );
189 clReleaseEvent(copyEvent);
190 clReleaseMemObject( streams[0] );
191 clReleaseMemObject( streams[1] );
192 free( (void *)int_output_ptr );
193 free( (void *)int_input_ptr );
194 return -1;
195 }
196
197 err = clGetEventProfilingInfo( copyEvent, CL_PROFILING_COMMAND_END, sizeof( cl_ulong ), &writeEnd, NULL );
198 if( err != CL_SUCCESS ){
199 print_error( err, "clGetEventProfilingInfo failed" );
200 clReleaseEvent(copyEvent);
201 clReleaseMemObject( streams[0] );
202 clReleaseMemObject( streams[1] );
203 free( (void *)int_output_ptr );
204 free( (void *)int_input_ptr );
205 return -1;
206 }
207
208 err = clEnqueueReadBuffer( queue, streams[1], true, 0, sizeof(cl_int)*num_elements, (void *)int_output_ptr, 0, NULL, NULL );
209 if( err != CL_SUCCESS ){
210 print_error( err, "clEnqueueReadBuffer failed" );
211 clReleaseEvent(copyEvent);
212 clReleaseMemObject( streams[0] );
213 clReleaseMemObject( streams[1] );
214 free( (void *)int_output_ptr );
215 free( (void *)int_input_ptr );
216 return -1;
217 }
218
219 if( verify_copy_array(int_input_ptr, int_output_ptr, num_elements) ){
220 log_error( "test failed\n" );
221 err = -1;
222 }
223 else{
224 log_info( "test passed\n" );
225 err = 0;
226 }
227
228 // cleanup
229 clReleaseEvent(copyEvent);
230 clReleaseMemObject( streams[0] );
231 clReleaseMemObject( streams[1] );
232 free( (void *)int_output_ptr );
233 free( (void *)int_input_ptr );
234
235 if (check_times(queueStart, submitStart, writeStart, writeEnd, device))
236 err = -1;
237
238 return err;
239
240 } // end copy_size()
241
242
copy_partial_size(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements,cl_uint srcStart,cl_uint dstStart,int size,MTdata d)243 static int copy_partial_size( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements, cl_uint srcStart, cl_uint dstStart, int size, MTdata d )
244 {
245 cl_mem streams[2];
246 cl_event copyEvent;
247 cl_ulong queueStart, submitStart, writeStart, writeEnd;
248 cl_int *inptr, *outptr;
249 int err = 0;
250 int i;
251
252 inptr = (cl_int *)malloc(sizeof(cl_int) * num_elements);
253 outptr = (cl_int *)malloc(sizeof(cl_int) * num_elements);
254
255 streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
256 sizeof(cl_int) * num_elements, NULL, &err);
257 if (!streams[0])
258 {
259 log_error("clCreateBuffer failed\n");
260 return -1;
261 }
262 streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
263 sizeof(cl_int) * num_elements, NULL, &err);
264 if (!streams[1])
265 {
266 log_error("clCreateBuffer failed\n");
267 return -1;
268 }
269
270 for (i=0; i<num_elements; i++){
271 inptr[i] = (int)genrand_int32(d);
272 outptr[i] = (int)get_random_float( -1.f, 1.f, d ); // seed with incorrect data
273 }
274
275 err = clEnqueueWriteBuffer(queue, streams[0], true, 0, sizeof(cl_int)*num_elements, (void *)inptr, 0, NULL, NULL);
276 if (err != CL_SUCCESS)
277 {
278 log_error("clWriteArray failed\n");
279 return -1;
280 }
281
282 err = clEnqueueCopyBuffer( queue, streams[0], streams[1], srcStart*sizeof(cl_int), dstStart*sizeof(cl_int),
283 sizeof(cl_int)*size, 0, NULL, ©Event );
284 if( err != CL_SUCCESS){
285 print_error( err, "clCopyArray failed" );
286 clReleaseMemObject( streams[0] );
287 clReleaseMemObject( streams[1] );
288 free( outptr );
289 free( inptr );
290 return -1;
291 }
292
293 // This synchronization point is needed in order to assume the data is valid.
294 // Getting profiling information is not a synchronization point.
295 err = clWaitForEvents( 1, ©Event );
296 if( err != CL_SUCCESS )
297 {
298 clReleaseEvent(copyEvent);
299 clReleaseMemObject( streams[0] );
300 clReleaseMemObject( streams[1] );
301 free( outptr );
302 free( inptr );
303 return -1;
304 }
305
306 // test profiling
307 while( ( err = clGetEventProfilingInfo( copyEvent, CL_PROFILING_COMMAND_QUEUED, sizeof( cl_ulong ), &queueStart, NULL ) ) ==
308 CL_PROFILING_INFO_NOT_AVAILABLE );
309 if( err != CL_SUCCESS ){
310 print_error( err, "clGetEventProfilingInfo failed" );
311 clReleaseEvent(copyEvent);
312 clReleaseMemObject( streams[0] );
313 clReleaseMemObject( streams[1] );
314 free( outptr );
315 free( inptr );
316 return -1;
317 }
318
319 while( ( err = clGetEventProfilingInfo( copyEvent, CL_PROFILING_COMMAND_SUBMIT, sizeof( cl_ulong ), &submitStart, NULL ) ) ==
320 CL_PROFILING_INFO_NOT_AVAILABLE );
321 if( err != CL_SUCCESS ){
322 print_error( err, "clGetEventProfilingInfo failed" );
323 clReleaseEvent(copyEvent);
324 clReleaseMemObject( streams[0] );
325 clReleaseMemObject( streams[1] );
326 free( outptr );
327 free( inptr );
328 return -1;
329 }
330
331
332 err = clGetEventProfilingInfo( copyEvent, CL_PROFILING_COMMAND_START, sizeof( cl_ulong ), &writeStart, NULL );
333 if( err != CL_SUCCESS ){
334 print_error( err, "clGetEventProfilingInfo failed" );
335 clReleaseEvent(copyEvent);
336 clReleaseMemObject( streams[0] );
337 clReleaseMemObject( streams[1] );
338 free( outptr );
339 free( inptr );
340 return -1;
341 }
342
343 err = clGetEventProfilingInfo( copyEvent, CL_PROFILING_COMMAND_END, sizeof( cl_ulong ), &writeEnd, NULL );
344 if( err != CL_SUCCESS ){
345 print_error( err, "clGetEventProfilingInfo failed" );
346 clReleaseEvent(copyEvent);
347 clReleaseMemObject( streams[0] );
348 clReleaseMemObject( streams[1] );
349 free( outptr );
350 free( inptr );
351 return -1;
352 }
353
354 err = clEnqueueReadBuffer( queue, streams[1], true, 0, sizeof(cl_int)*num_elements, (void *)outptr, 0, NULL, NULL );
355 if( err != CL_SUCCESS){
356 log_error("clReadVariableStream failed\n");
357 return -1;
358 }
359
360 if( verify_copy_array(inptr + srcStart, outptr + dstStart, size) ){
361 log_error("test failed\n");
362 err = -1;
363 }
364 else{
365 log_info("test passed\n");
366 err = 0;
367 }
368
369 // cleanup
370 clReleaseEvent(copyEvent);
371 clReleaseMemObject(streams[0]);
372 clReleaseMemObject(streams[1]);
373 free(outptr);
374 free(inptr);
375
376 if (check_times(queueStart, submitStart, writeStart, writeEnd, device))
377 err = -1;
378
379 return err;
380
381 } // end copy_partial_size()
382
383
test_copy_array(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)384 int test_copy_array( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
385 {
386 int i, err = 0;
387 int size;
388 MTdata d = init_genrand( gRandomSeed );
389
390 // test the preset size
391 log_info( "set size: %d: ", num_elements );
392 err = copy_size( device, context, queue, num_elements, d );
393
394 // now test random sizes
395 for( i = 0; i < 8; i++ ){
396 size = (int)get_random_float(2.f,131072.f, d);
397 log_info( "random size: %d: ", size );
398 err |= copy_size( device, context, queue, size, d );
399 }
400
401 free_mtdata(d);
402
403 return err;
404
405 } // end copy_array()
406
407
test_copy_partial_array(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)408 int test_copy_partial_array( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
409 {
410 int i, err = 0;
411 int size;
412 cl_uint srcStart, dstStart;
413 MTdata d = init_genrand( gRandomSeed );
414
415 // now test copy of partial sizes
416 for( i = 0; i < 8; i++ ){
417 srcStart = (cl_uint)get_random_float( 0.f, (float)(num_elements - 8), d );
418 size = (int)get_random_float( 8.f, (float)(num_elements - srcStart), d );
419 dstStart = (cl_uint)get_random_float( 0.f, (float)(num_elements - size), d );
420 log_info( "random partial copy from %d to %d, size: %d: ", (int)srcStart, (int)dstStart, size );
421 err |= copy_partial_size( device, context, queue, num_elements, srcStart, dstStart, size, d );
422 }
423
424 free_mtdata(d);
425 return err;
426 } // end copy_partial_array()
427
428
copy_image_size(cl_device_id device,cl_context context,cl_command_queue queue,size_t srcx,size_t srcy,size_t dstx,size_t dsty,size_t subw,size_t subh,MTdata d)429 static int copy_image_size( cl_device_id device, cl_context context,
430 cl_command_queue queue, size_t srcx, size_t srcy,
431 size_t dstx, size_t dsty, size_t subw, size_t subh,
432 MTdata d )
433 {
434 cl_mem memobjs[3];
435 cl_program program[1];
436 cl_image_format image_format_desc = { CL_RGBA, CL_UNORM_INT8 };
437 cl_event copyEvent;
438 cl_ulong queueStart, submitStart, writeStart, writeEnd;
439 void *inptr;
440 void *dst = NULL;
441 cl_kernel kernel[1];
442 size_t threads[2];
443 int err = 0;
444 cl_mem_flags flags;
445 unsigned int num_channels = 4;
446 size_t w = 256, h = 256;
447 size_t element_nbytes;
448 size_t num_bytes;
449 size_t channel_nbytes = sizeof( cl_char );
450
451
452 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
453
454 element_nbytes = channel_nbytes * num_channels;
455 num_bytes = w * h * element_nbytes;
456
457 threads[0] = (size_t)w;
458 threads[1] = (size_t)h;
459
460 inptr = (void *)generate_image( (int)num_bytes, d );
461 if( ! inptr ){
462 log_error("unable to allocate inptr at %d x %d\n", (int)w, (int)h );
463 return -1;
464 }
465
466 dst = malloc( num_bytes );
467 if( ! dst ){
468 free( (void *)inptr );
469 log_error("unable to allocate dst at %d x %d\n", (int)w, (int)h );
470 return -1;
471 }
472
473 // allocate the input image
474 flags = CL_MEM_READ_WRITE;
475 memobjs[0] = create_image_2d(context, flags, &image_format_desc, w, h, 0, NULL, &err);
476 if( memobjs[0] == (cl_mem)0 ) {
477 free( dst );
478 free( (void *)inptr );
479 log_error("unable to create Image2D\n");
480 return -1;
481 }
482
483 memobjs[1] =
484 clCreateBuffer(context, CL_MEM_READ_WRITE, num_bytes, NULL, &err);
485 if( memobjs[1] == (cl_mem)0 ) {
486 clReleaseMemObject(memobjs[0]);
487 free( dst );
488 free( (void *)inptr );
489 log_error("unable to create array\n");
490 return -1;
491 }
492
493 // allocate the input image
494 memobjs[2] = create_image_2d(context, flags, &image_format_desc, w, h, 0, NULL, &err);
495 if( memobjs[2] == (cl_mem)0 ) {
496 clReleaseMemObject(memobjs[0]);
497 clReleaseMemObject(memobjs[1]);
498 free( dst );
499 free( (void *)inptr );
500 log_error("unable to create Image2D\n");
501 return -1;
502 }
503
504 err = clEnqueueWriteBuffer( queue, memobjs[1], true, 0, num_bytes, inptr, 0, NULL, NULL );
505 if( err != CL_SUCCESS ){
506 log_error("clWriteArray failed\n");
507 return -1;
508 }
509
510 err = create_single_kernel_helper( context, &program[0], &kernel[0], 1, &write_kernel_code, "test_write" );
511 if( err ){
512 clReleaseMemObject( memobjs[0] );
513 clReleaseMemObject( memobjs[1] );
514 clReleaseMemObject( memobjs[2] );
515 free( dst );
516 free( inptr );
517 return -1;
518 }
519
520 err = clSetKernelArg( kernel[0], 0, sizeof( cl_mem ), (void *)&memobjs[1] );
521 err |= clSetKernelArg( kernel[0], 1, sizeof( cl_mem ), (void *)&memobjs[0] );
522 if (err != CL_SUCCESS){
523 log_error("clSetKernelArg failed\n");
524 clReleaseKernel( kernel[0] );
525 clReleaseProgram( program[0] );
526 clReleaseMemObject( memobjs[0] );
527 clReleaseMemObject( memobjs[1] );
528 clReleaseMemObject( memobjs[2] );
529 free( dst );
530 free( inptr );
531 return -1;
532 }
533
534 err = clEnqueueNDRangeKernel( queue, kernel[0], 2, NULL, threads, NULL, 0, NULL, NULL );
535
536 if (err != CL_SUCCESS){
537 print_error( err, "clEnqueueNDRangeKernel failed" );
538 clReleaseKernel( kernel[0] );
539 clReleaseProgram( program[0] );
540 clReleaseMemObject( memobjs[0] );
541 clReleaseMemObject( memobjs[1] );
542 clReleaseMemObject( memobjs[2] );
543 free( dst );
544 free( inptr );
545 return -1;
546 }
547
548 // now do the copy
549 size_t srcPt[3] = { srcx, srcy, 0 };
550 size_t destPt[3] = { dstx, dsty, 0 };
551 size_t region[3] = { subw, subh, 1 };
552 err = clEnqueueCopyImage( queue, memobjs[0], memobjs[2], srcPt, destPt, region, 0, NULL, ©Event );
553 if (err != CL_SUCCESS){
554 print_error( err, "clCopyImage failed" );
555 clReleaseKernel( kernel[0] );
556 clReleaseProgram( program[0] );
557 clReleaseMemObject( memobjs[0] );
558 clReleaseMemObject( memobjs[1] );
559 clReleaseMemObject( memobjs[2] );
560 free( dst );
561 free( inptr );
562 return -1;
563 }
564
565 // This synchronization point is needed in order to assume the data is valid.
566 // Getting profiling information is not a synchronization point.
567 err = clWaitForEvents( 1, ©Event );
568 if( err != CL_SUCCESS )
569 {
570 clReleaseEvent(copyEvent);
571 clReleaseKernel( kernel[0] );
572 clReleaseProgram( program[0] );
573 clReleaseMemObject( memobjs[0] );
574 clReleaseMemObject( memobjs[1] );
575 clReleaseMemObject( memobjs[2] );
576 free( dst );
577 free( inptr );
578 return -1;
579 }
580
581 // test profiling
582 while( ( err = clGetEventProfilingInfo( copyEvent, CL_PROFILING_COMMAND_QUEUED, sizeof( cl_ulong ), &queueStart, NULL ) ) ==
583 CL_PROFILING_INFO_NOT_AVAILABLE );
584 if( err != CL_SUCCESS ){
585 print_error( err, "clGetEventProfilingInfo failed" );
586 clReleaseEvent(copyEvent);
587 clReleaseKernel( kernel[0] );
588 clReleaseProgram( program[0] );
589 clReleaseMemObject( memobjs[0] );
590 clReleaseMemObject( memobjs[1] );
591 clReleaseMemObject( memobjs[2] );
592 free( dst );
593 free( inptr );
594 return -1;
595 }
596
597 while( ( err = clGetEventProfilingInfo( copyEvent, CL_PROFILING_COMMAND_SUBMIT, sizeof( cl_ulong ), &submitStart, NULL ) ) ==
598 CL_PROFILING_INFO_NOT_AVAILABLE );
599 if( err != CL_SUCCESS ){
600 print_error( err, "clGetEventProfilingInfo failed" );
601 clReleaseEvent(copyEvent);
602 clReleaseKernel( kernel[0] );
603 clReleaseProgram( program[0] );
604 clReleaseMemObject( memobjs[0] );
605 clReleaseMemObject( memobjs[1] );
606 clReleaseMemObject( memobjs[2] );
607 free( dst );
608 free( inptr );
609 return -1;
610 }
611
612 err = clGetEventProfilingInfo( copyEvent, CL_PROFILING_COMMAND_START, sizeof( cl_ulong ), &writeStart, NULL );
613 if( err != CL_SUCCESS ){
614 print_error( err, "clGetEventProfilingInfo failed" );
615 clReleaseEvent(copyEvent);
616 clReleaseKernel( kernel[0] );
617 clReleaseProgram( program[0] );
618 clReleaseMemObject( memobjs[0] );
619 clReleaseMemObject( memobjs[1] );
620 clReleaseMemObject( memobjs[2] );
621 free( dst );
622 free( inptr );
623 return -1;
624 }
625
626 err = clGetEventProfilingInfo( copyEvent, CL_PROFILING_COMMAND_END, sizeof( cl_ulong ), &writeEnd, NULL );
627 if( err != CL_SUCCESS ){
628 print_error( err, "clGetEventProfilingInfo failed" );
629 clReleaseEvent(copyEvent);
630 clReleaseKernel( kernel[0] );
631 clReleaseProgram( program[0] );
632 clReleaseMemObject( memobjs[0] );
633 clReleaseMemObject( memobjs[1] );
634 clReleaseMemObject( memobjs[2] );
635 free( dst );
636 free( inptr );
637 return -1;
638 }
639
640 size_t origin[3] = { 0, 0, 0 };
641 size_t region2[3] = { w, h, 1 };
642 err = clEnqueueReadImage( queue, memobjs[2], true, origin, region2, 0, 0, dst, 0, NULL, NULL );
643 if (err != CL_SUCCESS){
644 print_error( err, "clReadImage failed" );
645 clReleaseEvent(copyEvent);
646 clReleaseKernel( kernel[0] );
647 clReleaseProgram( program[0] );
648 clReleaseMemObject( memobjs[0] );
649 clReleaseMemObject( memobjs[1] );
650 clReleaseMemObject( memobjs[2] );
651 free( dst );
652 free( inptr );
653 return -1;
654 }
655
656 err = verify_subimage( (unsigned char *)inptr, (unsigned char *)dst, srcx, srcy,
657 dstx, dsty, subw, subh, w, 4 );
658 //err = verify_image( (unsigned char *)inptr, (unsigned char *)dst, w * h * 4 );
659 if( err ){
660 log_error( "Image failed to verify.\n " );
661 }
662 else{
663 log_info( "Image verified.\n" );
664 }
665
666 // cleanup
667 clReleaseEvent(copyEvent);
668 clReleaseKernel( kernel[0] );
669 clReleaseProgram( program[0] );
670 clReleaseMemObject( memobjs[0] );
671 clReleaseMemObject( memobjs[1] );
672 clReleaseMemObject( memobjs[2] );
673 free( dst );
674 free( inptr );
675
676 if (check_times(queueStart, submitStart, writeStart, writeEnd, device))
677 err = -1;
678
679 return err;
680
681 } // end copy_image_size()
682
683
test_copy_image(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)684 int test_copy_image( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
685 {
686 int err = 0;
687 int i;
688 size_t srcx, srcy, dstx, dsty, subw, subh;
689 MTdata d;
690
691 srcx = srcy = dstx = dsty = 0;
692 subw = subh = 256;
693
694 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
695
696 d = init_genrand( gRandomSeed );
697 err = copy_image_size( device, context, queue, srcx, srcy, dstx, dsty, subw, subh, d );
698 if( err ){
699 log_error( "testing copy image, full size\n" );
700 }
701 else{
702 log_info( "testing copy image, full size\n" );
703 }
704
705 // now test random sub images
706 srcx = srcy = 0;
707 subw = subh = 16;
708 dstx = dsty = 0;
709 err = copy_image_size( device, context, queue, srcx, srcy, dstx, dsty, subw, subh, d );
710 if( err ){
711 log_error( "test copy of subimage size %d,%d %d,%d %d x %d\n", (int)srcx, (int)srcy,
712 (int)dstx, (int)dsty, (int)subw, (int)subh );
713 }
714 else{
715 log_info( "test copy of subimage size %d,%d %d,%d %d x %d\n", (int)srcx, (int)srcy,
716 (int)dstx, (int)dsty, (int)subw, (int)subh );
717 }
718
719 srcx = srcy = 8;
720 subw = subh = 16;
721 dstx = dsty = 32;
722 err = copy_image_size( device, context, queue, srcx, srcy, dstx, dsty, subw, subh, d );
723 if( err ){
724 log_error( "test copy of subimage size %d,%d %d,%d %d x %d\n", (int)srcx, (int)srcy,
725 (int)dstx, (int)dsty, (int)subw, (int)subh );
726 }
727 else{
728 log_info( "test copy of subimage size %d,%d %d,%d %d x %d\n", (int)srcx, (int)srcy,
729 (int)dstx, (int)dsty, (int)subw, (int)subh );
730 }
731
732 for( i = 0; i < 16; i++ ) {
733 srcx = (size_t)get_random_float( 0.f, 248.f, d );
734 srcy = (size_t)get_random_float( 0.f, 248.f, d );
735 subw = (size_t)get_random_float( 8.f, (float)(256 - srcx), d );
736 subh = (size_t)get_random_float( 8.f, (float)(256 - srcy), d );
737 dstx = (size_t)get_random_float( 0.f, (float)(256 - subw), d );
738 dsty = (size_t)get_random_float( 0.f, (float)(256 - subh), d );
739 err = copy_image_size( device, context, queue, srcx, srcy, dstx, dsty, subw, subh, d );
740 if( err ){
741 log_error( "test copy of subimage size %d,%d %d,%d %d x %d\n", (int)srcx, (int)srcy,
742 (int)dstx, (int)dsty, (int)subw, (int)subh );
743 }
744 else{
745 log_info( "test copy of subimage size %d,%d %d,%d %d x %d\n", (int)srcx, (int)srcy,
746 (int)dstx, (int)dsty, (int)subw, (int)subh );
747 }
748 }
749
750 free_mtdata(d);
751
752 return err;
753
754 } // end copy_image()
755
756
test_copy_array_to_image(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)757 int test_copy_array_to_image( cl_device_id device, cl_context context, cl_command_queue queue, int num_elements )
758 {
759 cl_mem memobjs[3];
760 cl_image_format image_format_desc = { CL_RGBA, CL_UNORM_INT8 };
761 void *inptr;
762 void *dst;
763 int err;
764 cl_mem_flags flags;
765 unsigned int num_channels = (unsigned int)get_format_channel_count( &image_format_desc );
766 size_t w = 256, h = 256;
767 size_t element_nbytes;
768 size_t num_bytes;
769 size_t channel_nbytes = sizeof( cl_char );
770 MTdata d;
771
772 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
773
774 element_nbytes = channel_nbytes * num_channels;
775 num_bytes = w * h * element_nbytes;
776 d = init_genrand( gRandomSeed );
777 inptr = (void *)generate_image( (int)num_bytes, d );
778 free_mtdata(d); d = NULL;
779 if( ! inptr ){
780 log_error("unable to allocate inptr at %d x %d\n", (int)w, (int)h );
781 return -1;
782 }
783
784 dst = malloc( num_bytes );
785 if( ! dst ){
786 free( inptr );
787 log_error( " unable to allocate dst at %d x %d\n", (int)w, (int)h );
788 return -1;
789 }
790
791 // allocate the input image
792 flags = CL_MEM_READ_WRITE;
793 memobjs[0] = create_image_2d( context, flags, &image_format_desc, w, h, 0, NULL, &err );
794 if( memobjs[0] == (cl_mem)0 ){
795 free( dst );
796 free( inptr );
797 log_error( " unable to create Image2D\n" );
798 return -1;
799 }
800
801 memobjs[1] =
802 clCreateBuffer(context, CL_MEM_READ_WRITE,
803 channel_nbytes * num_channels * w * h, NULL, &err);
804 if( memobjs[1] == (cl_mem)0 ) {
805 clReleaseMemObject( memobjs[0] );
806 free( dst );
807 free( inptr );
808 log_error( " unable to create array: " );
809 return -1;
810 }
811
812 err = clEnqueueWriteBuffer( queue, memobjs[1], true, 0, num_bytes, (const void *)inptr, 0, NULL, NULL );
813 if( err != CL_SUCCESS ){
814 print_error( err, "clWriteArray failed" );
815 clReleaseMemObject( memobjs[1] );
816 clReleaseMemObject( memobjs[0] );
817 free( dst );
818 free( inptr );
819 return -1;
820 }
821
822 size_t origin[3] = { 0, 0, 0 };
823 size_t region[3] = { w, h, 1 };
824 err = clEnqueueCopyBufferToImage( queue, memobjs[1], memobjs[0], 0, origin, region, 0, NULL, NULL );
825 if( err != CL_SUCCESS ){
826 print_error( err, "clCopyArrayToImage failed" );
827 clReleaseMemObject( memobjs[1] );
828 clReleaseMemObject( memobjs[0] );
829 free( dst );
830 free( inptr );
831 return -1;
832 }
833
834 err = clEnqueueReadImage( queue, memobjs[0], true, origin, region, 0, 0, dst, 0, NULL, NULL );
835 if( err != CL_SUCCESS ){
836 print_error( err, "clReadImage failed" );
837 clReleaseMemObject( memobjs[1] );
838 clReleaseMemObject( memobjs[0] );
839 free( dst );
840 free( inptr );
841 return -1;
842 }
843
844 err = verify_subimage( (cl_uchar *)inptr, (cl_uchar *)dst, 0, 0, 0, 0, w, h, w, num_channels );
845 if( err ){
846 log_error( " test failed: " );
847 }
848 else{
849 log_info( " test passed: " );
850 }
851
852 // cleanup
853 clReleaseMemObject( memobjs[1] );
854 clReleaseMemObject( memobjs[0] );
855 free( dst );
856 free( inptr );
857
858 return err;
859
860 } // end copy_array_to_image()
861