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 int err;
417 int w = 64, h = 64;
418 cl_mem_flags flags;
419 size_t element_nbytes;
420 size_t num_bytes;
421 size_t channel_nbytes = sizeof( cl_uchar );
422 MTdata d;
423
424
425 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
426
427 if (readFloat)
428 channel_nbytes = sizeof( cl_float );
429
430 element_nbytes = channel_nbytes * get_format_channel_count( &image_format_desc );
431 num_bytes = w * h * element_nbytes;
432
433 threads[0] = (size_t)w;
434 threads[1] = (size_t)h;
435
436 d = init_genrand( gRandomSeed );
437 if( image_format_desc.image_channel_data_type == CL_SIGNED_INT8 )
438 inptr = (void *)generateSignedImage( w * h * 4, d );
439 else
440 inptr = (void *)generateImage( w * h * 4, d );
441 free_mtdata(d); d = NULL;
442 if( ! inptr ){
443 log_error("unable to allocate inptr at %d x %d\n", (int)w, (int)h );
444 return -1;
445 }
446
447 dst = malloc( num_bytes );
448 if( ! dst ){
449 free( (void *)inptr );
450 log_error("unable to allocate dst at %d x %d\n", (int)w, (int)h );
451 return -1;
452 }
453
454 // allocate the input and output image memory objects
455 flags = CL_MEM_READ_WRITE;
456 memobjs[0] = create_image_2d( context, flags, &image_format_desc, w, h, 0, NULL, &err );
457 if( memobjs[0] == (cl_mem)0 ){
458 free( dst );
459 free( (void *)inptr );
460 log_error("unable to create Image2D\n");
461 return -1;
462 }
463
464 memobjs[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
465 channel_nbytes * 4 * w * h, NULL, &err);
466 if( memobjs[1] == (cl_mem)0 ){
467 free( dst );
468 free( (void *)inptr );
469 clReleaseMemObject(memobjs[0]);
470 log_error("unable to create array\n");
471 return -1;
472 }
473
474 size_t origin[3] = { 0, 0, 0 };
475 size_t region[3] = { w, h, 1 };
476 err = clEnqueueWriteImage( queue, memobjs[0], false, origin, region, 0, 0, inptr, 0, NULL, &writeEvent );
477 if( err != CL_SUCCESS ){
478 clReleaseMemObject(memobjs[0]);
479 clReleaseMemObject(memobjs[1]);
480 free( dst );
481 free( inptr );
482 print_error(err, "clWriteImage failed");
483 return -1;
484 }
485
486 // This synchronization point is needed in order to assume the data is valid.
487 // Getting profiling information is not a synchronization point.
488 err = clWaitForEvents( 1, &writeEvent );
489 if( err != CL_SUCCESS )
490 {
491 print_error( err, "clWaitForEvents failed" );
492 clReleaseEvent(writeEvent);
493 clReleaseMemObject(memobjs[0]);
494 clReleaseMemObject(memobjs[1]);
495 free( dst );
496 free( inptr );
497 return -1;
498 }
499
500 // test profiling
501 while( ( err = clGetEventProfilingInfo( writeEvent, CL_PROFILING_COMMAND_QUEUED, sizeof( cl_ulong ), &queueStart, NULL ) ) ==
502 CL_PROFILING_INFO_NOT_AVAILABLE );
503 if( err != CL_SUCCESS ){
504 print_error( err, "clGetEventProfilingInfo failed" );
505 clReleaseEvent(writeEvent);
506 clReleaseMemObject(memobjs[0]);
507 clReleaseMemObject(memobjs[1]);
508 free( dst );
509 free( inptr );
510 return -1;
511 }
512
513 while( ( err = clGetEventProfilingInfo( writeEvent, CL_PROFILING_COMMAND_SUBMIT, sizeof( cl_ulong ), &submitStart, 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 err = clGetEventProfilingInfo( writeEvent, CL_PROFILING_COMMAND_START, sizeof( cl_ulong ), &writeStart, NULL );
526 if( err != CL_SUCCESS ){
527 print_error( err, "clGetEventProfilingInfo failed" );
528 clReleaseEvent(writeEvent);
529 clReleaseMemObject(memobjs[0]);
530 clReleaseMemObject(memobjs[1]);
531 free( dst );
532 free( inptr );
533 return -1;
534 }
535
536 err = clGetEventProfilingInfo( writeEvent, CL_PROFILING_COMMAND_END, sizeof( cl_ulong ), &writeEnd, NULL );
537 if( err != CL_SUCCESS ){
538 print_error( err, "clGetEventProfilingInfo failed" );
539 clReleaseEvent(writeEvent);
540 clReleaseMemObject(memobjs[0]);
541 clReleaseMemObject(memobjs[1]);
542 free( dst );
543 free( inptr );
544 return -1;
545 }
546
547 err = create_single_kernel_helper( context, &program[0], &kernel[0], 1, &code, name );
548 if( err ){
549 log_error( "Unable to create program and kernel\n" );
550 clReleaseEvent(writeEvent);
551 clReleaseMemObject(memobjs[0]);
552 clReleaseMemObject(memobjs[1]);
553 free( dst );
554 free( inptr );
555 return -1;
556 }
557
558 err = clSetKernelArg( kernel[0], 0, sizeof( cl_mem ), (void *)&memobjs[0] );
559 err |= clSetKernelArg( kernel[0], 1, sizeof( cl_mem ), (void *)&memobjs[1] );
560 if( err != CL_SUCCESS ){
561 log_error( "clSetKernelArg failed\n" );
562 clReleaseEvent(writeEvent);
563 clReleaseKernel( kernel[0] );
564 clReleaseProgram( program[0] );
565 clReleaseMemObject(memobjs[0]);
566 clReleaseMemObject(memobjs[1]);
567 free( dst );
568 free( inptr );
569 return -1;
570 }
571
572 err = clEnqueueNDRangeKernel(queue, kernel[0], 2, NULL, threads, NULL, 0, NULL, NULL );
573
574 if( err != CL_SUCCESS ){
575 print_error( err, "clEnqueueNDRangeKernel failed" );
576 clReleaseEvent(writeEvent);
577 clReleaseKernel( kernel[0] );
578 clReleaseProgram( program[0] );
579 clReleaseMemObject(memobjs[0]);
580 clReleaseMemObject(memobjs[1]);
581 free( dst );
582 free( inptr );
583 return -1;
584 }
585
586 err = clEnqueueReadBuffer( queue, memobjs[1], true, 0, num_bytes, dst, 0, NULL, NULL );
587 if( err != CL_SUCCESS ){
588 print_error( err, "clEnqueueReadBuffer failed" );
589 clReleaseEvent(writeEvent);
590 clReleaseKernel( kernel[0] );
591 clReleaseProgram( program[0] );
592 clReleaseMemObject(memobjs[0]);
593 clReleaseMemObject(memobjs[1]);
594 free( dst );
595 free( inptr );
596 return -1;
597 }
598
599 if ( readFloat )
600 {
601 refptr = prepareReference( (cl_uchar *)inptr, w, h );
602 if ( refptr )
603 {
604 err = verifyImageFloat( refptr, (cl_float *)dst, w, h );
605 free ( refptr );
606 }
607 else
608 err = -1;
609 }
610 else
611 err = verifyImage( (cl_uchar *)inptr, (cl_uchar *)dst, w, h );
612
613 if( err )
614 {
615 log_error( "Image failed to verify.\n" );
616 }
617 else
618 {
619 log_info( "Image verified.\n" );
620 }
621
622 // cleanup
623 clReleaseEvent(writeEvent);
624 clReleaseKernel( kernel[0] );
625 clReleaseProgram( program[0] );
626 clReleaseMemObject(memobjs[0]);
627 clReleaseMemObject(memobjs[1]);
628 free( dst );
629 free( inptr );
630
631 if (check_times(queueStart, submitStart, writeStart, writeEnd, device))
632 err = -1;
633
634 return err;
635
636 } // end write_image()
637
638
test_write_image_float(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)639 int test_write_image_float( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
640 {
641 cl_image_format image_format_desc = { CL_RGBA, CL_UNORM_INT8 };
642 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
643 // 0 to 255 for unsigned image data
644 return write_image( device, context, queue, numElements, readKernelCode[0], readKernelName[0], image_format_desc, 1 );
645
646 }
647
648
test_write_image_char(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)649 int test_write_image_char( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
650 {
651 cl_image_format image_format_desc = { CL_RGBA, CL_SIGNED_INT8 };
652 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
653 // -128 to 127 for signed iamge data
654 return write_image( device, context, queue, numElements, readKernelCode[1], readKernelName[1], image_format_desc, 0 );
655
656 }
657
658
test_write_image_uchar(cl_device_id device,cl_context context,cl_command_queue queue,int numElements)659 int test_write_image_uchar( cl_device_id device, cl_context context, cl_command_queue queue, int numElements )
660 {
661 cl_image_format image_format_desc = { CL_RGBA, CL_UNSIGNED_INT8 };
662 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
663 // 0 to 255 for unsigned image data
664 return write_image( device, context, queue, numElements, readKernelCode[2], readKernelName[2], image_format_desc, 0 );
665
666 }
667
668
669