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
27 //--- the code for the kernel executables
28 static const char *readKernelCode[] = {
29 "__kernel void testReadf(read_only image2d_t srcimg, __global float4 *dst)\n"
30 "{\n"
31 " int tid_x = get_global_id(0);\n"
32 " int tid_y = get_global_id(1);\n"
33 " int indx = tid_y * get_image_width(srcimg) + tid_x;\n"
34 " float4 color;\n"
35 "\n"
36 " const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE;\n"
37 " color = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y));\n"
38 " dst[indx].x = color.x;\n"
39 " dst[indx].y = color.y;\n"
40 " dst[indx].z = color.z;\n"
41 " dst[indx].w = color.w;\n"
42 "\n"
43 "}\n",
44
45 "__kernel void testReadi(read_only image2d_t srcimg, __global uchar4 *dst)\n"
46 "{\n"
47 " int tid_x = get_global_id(0);\n"
48 " int tid_y = get_global_id(1);\n"
49 " int indx = tid_y * get_image_width(srcimg) + tid_x;\n"
50 " int4 color;\n"
51 "\n"
52 " const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE;\n"
53 " color = read_imagei(srcimg, sampler, (int2)(tid_x, tid_y));\n"
54 " uchar4 dst_write;\n"
55 " dst_write.x = (uchar)color.x;\n"
56 " dst_write.y = (uchar)color.y;\n"
57 " dst_write.z = (uchar)color.z;\n"
58 " dst_write.w = (uchar)color.w;\n"
59 " dst[indx] = dst_write;\n"
60 "\n"
61 "}\n",
62
63 "__kernel void testReadui(read_only image2d_t srcimg, __global uchar4 *dst)\n"
64 "{\n"
65 " int tid_x = get_global_id(0);\n"
66 " int tid_y = get_global_id(1);\n"
67 " int indx = tid_y * get_image_width(srcimg) + tid_x;\n"
68 " uint4 color;\n"
69 "\n"
70 " const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE;\n"
71 " color = read_imageui(srcimg, sampler, (int2)(tid_x, tid_y));\n"
72 " uchar4 dst_write;\n"
73 " dst_write.x = (uchar)color.x;\n"
74 " dst_write.y = (uchar)color.y;\n"
75 " dst_write.z = (uchar)color.z;\n"
76 " dst_write.w = (uchar)color.w;\n"
77 " dst[indx] = dst_write;\n"
78 "\n"
79 "}\n",
80
81 "__kernel void testWritef(__global uchar *src, write_only image2d_t dstimg)\n"
82 "{\n"
83 " int tid_x = get_global_id(0);\n"
84 " int tid_y = get_global_id(1);\n"
85 " int indx = tid_y * get_image_width(dstimg) + tid_x;\n"
86 " float4 color;\n"
87 "\n"
88 " indx *= 4;\n"
89 " color = (float4)((float)src[indx+0], (float)src[indx+1], (float)src[indx+2], (float)src[indx+3]);\n"
90 " color /= (float4)(255.f, 255.f, 255.f, 255.f);\n"
91 " write_imagef(dstimg, (int2)(tid_x, tid_y), color);\n"
92 "\n"
93 "}\n",
94
95 "__kernel void testWritei(__global char *src, write_only image2d_t dstimg)\n"
96 "{\n"
97 " int tid_x = get_global_id(0);\n"
98 " int tid_y = get_global_id(1);\n"
99 " int indx = tid_y * get_image_width(dstimg) + tid_x;\n"
100 " int4 color;\n"
101 "\n"
102 " indx *= 4;\n"
103 " color.x = (int)src[indx+0];\n"
104 " color.y = (int)src[indx+1];\n"
105 " color.z = (int)src[indx+2];\n"
106 " color.w = (int)src[indx+3];\n"
107 " write_imagei(dstimg, (int2)(tid_x, tid_y), color);\n"
108 "\n"
109 "}\n",
110
111 "__kernel void testWriteui(__global uchar *src, write_only image2d_t dstimg)\n"
112 "{\n"
113 " int tid_x = get_global_id(0);\n"
114 " int tid_y = get_global_id(1);\n"
115 " int indx = tid_y * get_image_width(dstimg) + tid_x;\n"
116 " uint4 color;\n"
117 "\n"
118 " indx *= 4;\n"
119 " color.x = (uint)src[indx+0];\n"
120 " color.y = (uint)src[indx+1];\n"
121 " color.z = (uint)src[indx+2];\n"
122 " color.w = (uint)src[indx+3];\n"
123 " write_imageui(dstimg, (int2)(tid_x, tid_y), color);\n"
124 "\n"
125 "}\n",
126
127 "__kernel void testReadWriteff(read_only image2d_t srcimg, write_only image2d_t dstimg)\n"
128 "{\n"
129 " int tid_x = get_global_id(0);\n"
130 " int tid_y = get_global_id(1);\n"
131 " float4 color;\n"
132 "\n"
133 " color = read_imagef(srcimg, CLK_DEFAULT_SAMPLER, (int2)(tid_x, tid_y));\n"
134 " write_imagef(dstimg, (int2)(tid_x, tid_y), color);\n"
135 "\n"
136 "}\n",
137
138 "__kernel void testReadWriteii(read_only image2d_t srcimg, write_only image2d_t dstimg)\n"
139 "{\n"
140 " int tid_x = get_global_id(0);\n"
141 " int tid_y = get_global_id(1);\n"
142 " int4 color;\n"
143 "\n"
144 " color = read_imagei(srcimg, CLK_DEFAULT_SAMPLER, (int2)(tid_x, tid_y));\n"
145 " write_imagei(dstimg, (int2)(tid_x, tid_y), color);\n"
146 "\n"
147 "}\n",
148
149 "__kernel void testReadWriteuiui(read_only image2d_t srcimg, write_only image2d_t dstimg)\n"
150 "{\n"
151 " int tid_x = get_global_id(0);\n"
152 " int tid_y = get_global_id(1);\n"
153 " uint4 color;\n"
154 "\n"
155 " color = read_imageui(srcimg, CLK_DEFAULT_SAMPLER, (int2)(tid_x, tid_y));\n"
156 " write_imageui(dstimg, (int2)(tid_x, tid_y), color);\n"
157 "\n"
158 "}\n",
159
160 "__kernel void testReadWritefi(read_only image2d_t srcimg, write_only image2d_t dstimg)\n"
161 "{\n"
162 " int tid_x = get_global_id(0);\n"
163 " int tid_y = get_global_id(1);\n"
164 " float4 colorf;\n"
165 " int4 colori;\n"
166 "\n"
167 " const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE;\n"
168 " colorf = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y));\n"
169 // since we are going from unsigned to signed, be sure to convert
170 // values greater 0.5 to negative values
171 " if( colorf.x >= 0.5f )\n"
172 " colori.x = (int)( ( colorf.x - 1.f ) * 255.f );\n"
173 " else\n"
174 " colori.x = (int)( colorf.x * 255.f );\n"
175 " if( colorf.y >= 0.5f )\n"
176 " colori.y = (int)( ( colorf.y - 1.f ) * 255.f );\n"
177 " else\n"
178 " colori.y = (int)( colorf.y * 255.f );\n"
179 " if( colorf.z >= 0.5f )\n"
180 " colori.z = (int)( ( colorf.z - 1.f ) * 255.f );\n"
181 " else\n"
182 " colori.z = (int)( colorf.z * 255.f );\n"
183 " if( colorf.w >= 0.5f )\n"
184 " colori.w = (int)( ( colorf.w - 1.f ) * 255.f );\n"
185 " else\n"
186 " colori.w = (int)( colorf.w * 255.f );\n"
187 " write_imagei(dstimg, (int2)(tid_x, tid_y), colori);\n"
188 "\n"
189 "}\n",
190
191 "__kernel void testReadWritefui(read_only image2d_t srcimg, write_only image2d_t dstimg)\n"
192 "{\n"
193 " int tid_x = get_global_id(0);\n"
194 " int tid_y = get_global_id(1);\n"
195 " float4 colorf;\n"
196 " uint4 colorui;\n"
197 "\n"
198 " const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE;\n"
199 " colorf = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y));\n"
200 " colorui.x = (uint)( colorf.x * 255.f );\n"
201 " colorui.y = (uint)( colorf.y * 255.f );\n"
202 " colorui.z = (uint)( colorf.z * 255.f );\n"
203 " colorui.w = (uint)( colorf.w * 255.f );\n"
204 " write_imageui(dstimg, (int2)(tid_x, tid_y), colorui);\n"
205 "\n"
206 "}\n",
207
208 "__kernel void testReadWriteif(read_only image2d_t srcimg, write_only image2d_t dstimg)\n"
209 "{\n"
210 " int tid_x = get_global_id(0);\n"
211 " int tid_y = get_global_id(1);\n"
212 " int4 colori;\n"
213 " float4 colorf;\n"
214 "\n"
215 // since we are going from signed to unsigned, we need to adjust the rgba values from
216 // from the signed image to add 256 to the signed image values less than 0.
217 " const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE;\n"
218 " colori = read_imagei(srcimg, sampler, (int2)(tid_x, tid_y));\n"
219 " if( colori.x < 0 )\n"
220 " colorf.x = ( (float)colori.x + 256.f ) / 255.f;\n"
221 " else\n"
222 " colorf.x = (float)colori.x / 255.f;\n"
223 " if( colori.y < 0 )\n"
224 " colorf.y = ( (float)colori.y + 256.f ) / 255.f;\n"
225 " else\n"
226 " colorf.y = (float)colori.y / 255.f;\n"
227 " if( colori.z < 0 )\n"
228 " colorf.z = ( (float)colori.z + 256.f ) / 255.f;\n"
229 " else\n"
230 " colorf.z = (float)colori.z / 255.f;\n"
231 " if( colori.w < 0 )\n"
232 " colorf.w = ( (float)colori.w + 256.f ) / 255.f;\n"
233 " else\n"
234 " colorf.w = (float)colori.w / 255.f;\n"
235 " write_imagef(dstimg, (int2)(tid_x, tid_y), colorf);\n"
236 "\n"
237 "}\n",
238
239 "__kernel void testReadWriteiui(read_only image2d_t srcimg, write_only image2d_t dstimg)\n"
240 "{\n"
241 " int tid_x = get_global_id(0);\n"
242 " int tid_y = get_global_id(1);\n"
243 " int4 colori;\n"
244 " uint4 colorui;\n"
245 "\n"
246 " const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE;\n"
247 " colori = read_imagei(srcimg, sampler, (int2)(tid_x, tid_y));\n"
248 // since we are going from signed to unsigned, we need to adjust the rgba values from
249 // from the signed image to add 256 to the signed image values less than 0.
250 " if( colori.x < 0 )\n"
251 " colorui.x = (uint)( colori.x + 256 );\n"
252 " else\n"
253 " colorui.x = (uint)colori.x;\n"
254 " if( colori.y < 0 )\n"
255 " colorui.y = (uint)( colori.y + 256 );\n"
256 " else\n"
257 " colorui.y = (uint)colori.y;\n"
258 " if( colori.z < 0 )\n"
259 " colorui.z = (uint)( colori.z + 256 );\n"
260 " else\n"
261 " colorui.z = (uint)colori.z;\n"
262 " if( colori.w < 0 )\n"
263 " colorui.w = (uint)( colori.w + 256 );\n"
264 " else\n"
265 " colorui.w = (uint)colori.w;\n"
266 " write_imageui(dstimg, (int2)(tid_x, tid_y), colorui);\n"
267 "\n"
268 "}\n",
269
270 "__kernel void testReadWriteuif(read_only image2d_t srcimg, write_only image2d_t dstimg)\n"
271 "{\n"
272 " int tid_x = get_global_id(0);\n"
273 " int tid_y = get_global_id(1);\n"
274 " uint4 colorui;\n"
275 " float4 colorf;\n"
276 "\n"
277 " const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE;\n"
278 " colorui = read_imageui(srcimg, sampler, (int2)(tid_x, tid_y));\n"
279 " colorf.x = (float)colorui.x / 255.f;\n"
280 " colorf.y = (float)colorui.y / 255.f;\n"
281 " colorf.z = (float)colorui.z / 255.f;\n"
282 " colorf.w = (float)colorui.w / 255.f;\n"
283 " write_imagef(dstimg, (int2)(tid_x, tid_y), colorf);\n"
284 "\n"
285 "}\n",
286
287 "__kernel void testReadWriteuii(read_only image2d_t srcimg, write_only image2d_t dstimg)\n"
288 "{\n"
289 " int tid_x = get_global_id(0);\n"
290 " int tid_y = get_global_id(1);\n"
291 " uint4 colorui;\n"
292 " int4 colori;\n"
293 "\n"
294 " const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST | CLK_NORMALIZED_COORDS_FALSE;\n"
295 " colorui = read_imageui(srcimg, sampler, (int2)(tid_x, tid_y));\n"
296 // since we are going from unsigned to signed, be sure to convert
297 // values greater 0.5 to negative values
298 " if( colorui.x >= 128U )\n"
299 " colori.x = (int)colorui.x - 256;\n"
300 " else\n"
301 " colori.x = (int)colorui.x;\n"
302 " if( colorui.y >= 128U )\n"
303 " colori.y = (int)colorui.y - 256;\n"
304 " else\n"
305 " colori.y = (int)colorui.y;\n"
306 " if( colorui.z >= 128U )\n"
307 " colori.z = (int)colorui.z - 256;\n"
308 " else\n"
309 " colori.z = (int)colorui.z;\n"
310 " if( colorui.w >= 128U )\n"
311 " colori.w = (int)colorui.w - 256;\n"
312 " else\n"
313 " colori.w = (int)colorui.w;\n"
314 " write_imagei(dstimg, (int2)(tid_x, tid_y), colori);\n"
315 "\n"
316 "}\n" };
317
318 static const char *readKernelName[] = { "testReadf", "testReadi", "testReadui", "testWritef", "testWritei", "testWriteui",
319 "testReadWriteff", "testReadWriteii", "testReadWriteuiui", "testReadWritefi",
320 "testReadWritefui", "testReadWriteif", "testReadWriteiui", "testReadWriteuif",
321 "testReadWriteuii" };
322
323
generateImage(int n,MTdata d)324 static cl_uchar *generateImage( int n, MTdata d )
325 {
326 cl_uchar *ptr = (cl_uchar *)malloc( n * sizeof( cl_uchar ) );
327 int i;
328
329 for( i = 0; i < n; i++ ){
330 ptr[i] = (cl_uchar)genrand_int32(d);
331 }
332
333 return ptr;
334
335 }
336
337
generateSignedImage(int n,MTdata d)338 static char *generateSignedImage( int n, MTdata d )
339 {
340 char *ptr = (char *)malloc( n * sizeof( char ) );
341 int i;
342
343 for( i = 0; i < n; i++ ){
344 ptr[i] = (char)genrand_int32(d);
345 }
346
347 return ptr;
348
349 }
350
351
verifyImage(cl_uchar * image,cl_uchar * outptr,int w,int h)352 static int verifyImage( cl_uchar *image, cl_uchar *outptr, int w, int h )
353 {
354 int i;
355
356 for( i = 0; i < w * h * 4; i++ ){
357 if( outptr[i] != image[i] ){
358 log_error("Image verification failed at offset %d. Actual value=%d, expected value=%d\n", i, outptr[i], image[i]);
359 return -1;
360 }
361 }
362
363 return 0;
364 }
365
verifyImageFloat(cl_double * refptr,cl_float * outptr,int w,int h)366 static int verifyImageFloat ( cl_double *refptr, cl_float *outptr, int w, int h )
367 {
368 int i;
369
370 for (i=0; i<w*h*4; i++)
371 {
372 if (outptr[i] != (float)refptr[i])
373 {
374 float ulps = Ulp_Error( outptr[i], refptr[i]);
375
376 if(! (fabsf(ulps) < 1.5f) )
377 {
378 log_error( "ERROR: Data sample %d does not validate! Expected (%a), got (%a), ulp %f\n",
379 (int)i, refptr[i], outptr[ i ], ulps );
380 return -1;
381 }
382 }
383 }
384
385 return 0;
386 }
387
prepareReference(cl_uchar * inptr,int w,int h)388 static double *prepareReference( cl_uchar *inptr, int w, int h)
389 {
390 int i;
391 double *refptr = (double *)malloc( w * h * 4*sizeof( double ) );
392 if ( !refptr )
393 {
394 log_error( "Unable to allocate refptr at %d x %d\n", (int)w, (int)h );
395 return 0;
396 }
397 for( i = 0; i < w * h * 4; i++ ) {
398 refptr[i] = ((double)inptr[i])/255;
399 }
400 return refptr;
401 }
402
403 //----- the test functions
write_image(cl_device_id device,cl_context context,cl_command_queue queue,int numElements,const char * code,const char * name,cl_image_format image_format_desc,int readFloat)404 int write_image( cl_device_id device, cl_context context, cl_command_queue queue, int numElements, const char *code,
405 const char *name, cl_image_format image_format_desc, int readFloat )
406 {
407 cl_mem memobjs[2];
408 cl_program program[1];
409 void *inptr;
410 double *refptr = NULL;
411 void *dst = NULL;
412 cl_kernel kernel[1];
413 cl_event writeEvent;
414 cl_ulong queueStart, submitStart, writeStart, writeEnd;
415 size_t threads[2];
416 #ifdef USE_LOCAL_THREADS
417 size_t localThreads[2];
418 #endif
419 int err;
420 int w = 64, h = 64;
421 cl_mem_flags flags;
422 size_t element_nbytes;
423 size_t num_bytes;
424 size_t channel_nbytes = sizeof( cl_uchar );
425 MTdata d;
426
427
428 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
429
430 if (readFloat)
431 channel_nbytes = sizeof( cl_float );
432
433 element_nbytes = channel_nbytes * get_format_channel_count( &image_format_desc );
434 num_bytes = w * h * element_nbytes;
435
436 threads[0] = (size_t)w;
437 threads[1] = (size_t)h;
438
439 #ifdef USE_LOCAL_THREADS
440 err = clGetDeviceConfigInfo( id, CL_DEVICE_MAX_THREAD_GROUP_SIZE, localThreads, sizeof( unsigned int ), NULL );
441 test_error( err, "Unable to get thread group max size" );
442 localThreads[1] = localThreads[0];
443 if( localThreads[0] > threads[0] )
444 localThreads[0] = threads[0];
445 if( localThreads[1] > threads[1] )
446 localThreads[1] = threads[1];
447 #endif
448
449 d = init_genrand( gRandomSeed );
450 if( image_format_desc.image_channel_data_type == CL_SIGNED_INT8 )
451 inptr = (void *)generateSignedImage( w * h * 4, d );
452 else
453 inptr = (void *)generateImage( w * h * 4, d );
454 free_mtdata(d); d = NULL;
455 if( ! inptr ){
456 log_error("unable to allocate inptr at %d x %d\n", (int)w, (int)h );
457 return -1;
458 }
459
460 dst = malloc( num_bytes );
461 if( ! dst ){
462 free( (void *)inptr );
463 log_error("unable to allocate dst at %d x %d\n", (int)w, (int)h );
464 return -1;
465 }
466
467 // allocate the input and output image memory objects
468 flags = (cl_mem_flags)(CL_MEM_READ_WRITE);
469 memobjs[0] = create_image_2d( context, flags, &image_format_desc, w, h, 0, NULL, &err );
470 if( memobjs[0] == (cl_mem)0 ){
471 free( dst );
472 free( (void *)inptr );
473 log_error("unable to create Image2D\n");
474 return -1;
475 }
476
477 memobjs[1] = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE), channel_nbytes * 4 * w * h, NULL, &err );
478 if( memobjs[1] == (cl_mem)0 ){
479 free( dst );
480 free( (void *)inptr );
481 clReleaseMemObject(memobjs[0]);
482 log_error("unable to create array\n");
483 return -1;
484 }
485
486 size_t origin[3] = { 0, 0, 0 };
487 size_t region[3] = { w, h, 1 };
488 err = clEnqueueWriteImage( queue, memobjs[0], false, origin, region, 0, 0, inptr, 0, NULL, &writeEvent );
489 if( err != CL_SUCCESS ){
490 clReleaseMemObject(memobjs[0]);
491 clReleaseMemObject(memobjs[1]);
492 free( dst );
493 free( inptr );
494 print_error(err, "clWriteImage failed");
495 return -1;
496 }
497
498 // This synchronization point is needed in order to assume the data is valid.
499 // Getting profiling information is not a synchronization point.
500 err = clWaitForEvents( 1, &writeEvent );
501 if( err != CL_SUCCESS )
502 {
503 print_error( err, "clWaitForEvents failed" );
504 clReleaseEvent(writeEvent);
505 clReleaseMemObject(memobjs[0]);
506 clReleaseMemObject(memobjs[1]);
507 free( dst );
508 free( inptr );
509 return -1;
510 }
511
512 // test profiling
513 while( ( err = clGetEventProfilingInfo( writeEvent, CL_PROFILING_COMMAND_QUEUED, sizeof( cl_ulong ), &queueStart, NULL ) ) ==
514 CL_PROFILING_INFO_NOT_AVAILABLE );
515 if( err != CL_SUCCESS ){
516 print_error( err, "clGetEventProfilingInfo failed" );
517 clReleaseEvent(writeEvent);
518 clReleaseMemObject(memobjs[0]);
519 clReleaseMemObject(memobjs[1]);
520 free( dst );
521 free( inptr );
522 return -1;
523 }
524
525 while( ( err = clGetEventProfilingInfo( writeEvent, CL_PROFILING_COMMAND_SUBMIT, sizeof( cl_ulong ), &submitStart, NULL ) ) ==
526 CL_PROFILING_INFO_NOT_AVAILABLE );
527 if( err != CL_SUCCESS ){
528 print_error( err, "clGetEventProfilingInfo failed" );
529 clReleaseEvent(writeEvent);
530 clReleaseMemObject(memobjs[0]);
531 clReleaseMemObject(memobjs[1]);
532 free( dst );
533 free( inptr );
534 return -1;
535 }
536
537 err = clGetEventProfilingInfo( writeEvent, CL_PROFILING_COMMAND_START, sizeof( cl_ulong ), &writeStart, NULL );
538 if( err != CL_SUCCESS ){
539 print_error( err, "clGetEventProfilingInfo failed" );
540 clReleaseEvent(writeEvent);
541 clReleaseMemObject(memobjs[0]);
542 clReleaseMemObject(memobjs[1]);
543 free( dst );
544 free( inptr );
545 return -1;
546 }
547
548 err = clGetEventProfilingInfo( writeEvent, CL_PROFILING_COMMAND_END, sizeof( cl_ulong ), &writeEnd, NULL );
549 if( err != CL_SUCCESS ){
550 print_error( err, "clGetEventProfilingInfo failed" );
551 clReleaseEvent(writeEvent);
552 clReleaseMemObject(memobjs[0]);
553 clReleaseMemObject(memobjs[1]);
554 free( dst );
555 free( inptr );
556 return -1;
557 }
558
559 err = create_single_kernel_helper( context, &program[0], &kernel[0], 1, &code, name );
560 if( err ){
561 log_error( "Unable to create program and kernel\n" );
562 clReleaseEvent(writeEvent);
563 clReleaseMemObject(memobjs[0]);
564 clReleaseMemObject(memobjs[1]);
565 free( dst );
566 free( inptr );
567 return -1;
568 }
569
570 err = clSetKernelArg( kernel[0], 0, sizeof( cl_mem ), (void *)&memobjs[0] );
571 err |= clSetKernelArg( kernel[0], 1, sizeof( cl_mem ), (void *)&memobjs[1] );
572 if( err != CL_SUCCESS ){
573 log_error( "clSetKernelArg failed\n" );
574 clReleaseEvent(writeEvent);
575 clReleaseKernel( kernel[0] );
576 clReleaseProgram( program[0] );
577 clReleaseMemObject(memobjs[0]);
578 clReleaseMemObject(memobjs[1]);
579 free( dst );
580 free( inptr );
581 return -1;
582 }
583
584 #ifdef USE_LOCAL_THREADS
585 err = clEnqueueNDRangeKernel(queue, kernel[0], 2, NULL, threads, localThreads, 0, NULL, NULL );
586 #else
587 err = clEnqueueNDRangeKernel(queue, kernel[0], 2, NULL, threads, NULL, 0, NULL, NULL );
588 #endif
589 if( err != CL_SUCCESS ){
590 print_error( err, "clEnqueueNDRangeKernel failed" );
591 clReleaseEvent(writeEvent);
592 clReleaseKernel( kernel[0] );
593 clReleaseProgram( program[0] );
594 clReleaseMemObject(memobjs[0]);
595 clReleaseMemObject(memobjs[1]);
596 free( dst );
597 free( inptr );
598 return -1;
599 }
600
601 err = clEnqueueReadBuffer( queue, memobjs[1], true, 0, num_bytes, dst, 0, NULL, NULL );
602 if( err != CL_SUCCESS ){
603 print_error( err, "clEnqueueReadBuffer failed" );
604 clReleaseEvent(writeEvent);
605 clReleaseKernel( kernel[0] );
606 clReleaseProgram( program[0] );
607 clReleaseMemObject(memobjs[0]);
608 clReleaseMemObject(memobjs[1]);
609 free( dst );
610 free( inptr );
611 return -1;
612 }
613
614 if ( readFloat )
615 {
616 refptr = prepareReference( (cl_uchar *)inptr, w, h );
617 if ( refptr )
618 {
619 err = verifyImageFloat( refptr, (cl_float *)dst, w, h );
620 free ( refptr );
621 }
622 else
623 err = -1;
624 }
625 else
626 err = verifyImage( (cl_uchar *)inptr, (cl_uchar *)dst, w, h );
627
628 if( err )
629 {
630 log_error( "Image failed to verify.\n" );
631 }
632 else
633 {
634 log_info( "Image verified.\n" );
635 }
636
637 // cleanup
638 clReleaseEvent(writeEvent);
639 clReleaseKernel( kernel[0] );
640 clReleaseProgram( program[0] );
641 clReleaseMemObject(memobjs[0]);
642 clReleaseMemObject(memobjs[1]);
643 free( dst );
644 free( inptr );
645
646 if (check_times(queueStart, submitStart, writeStart, writeEnd, device))
647 err = -1;
648
649 return err;
650
651 } // end write_image()
652
653
test_write_image_float(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)654 int test_write_image_float( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
655 {
656 cl_image_format image_format_desc = { CL_RGBA, CL_UNORM_INT8 };
657 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
658 // 0 to 255 for unsigned image data
659 return write_image( device, context, queue, numElements, readKernelCode[0], readKernelName[0], image_format_desc, 1 );
660
661 }
662
663
test_write_image_char(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)664 int test_write_image_char( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
665 {
666 cl_image_format image_format_desc = { CL_RGBA, CL_SIGNED_INT8 };
667 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
668 // -128 to 127 for signed iamge data
669 return write_image( device, context, queue, numElements, readKernelCode[1], readKernelName[1], image_format_desc, 0 );
670
671 }
672
673
test_write_image_uchar(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)674 int test_write_image_uchar( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
675 {
676 cl_image_format image_format_desc = { CL_RGBA, CL_UNSIGNED_INT8 };
677 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
678 // 0 to 255 for unsigned image data
679 return write_image( device, context, queue, numElements, readKernelCode[2], readKernelName[2], image_format_desc, 0 );
680
681 }
682
683
684