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
25 #include "procs.h"
26
27 static const char *image_to_image_kernel_integer_coord_code =
28 "\n"
29 "__kernel void image_to_image_copy(read_only image2d_t srcimg, write_only image2d_t dstimg, sampler_t sampler)\n"
30 "{\n"
31 " int tid_x = get_global_id(0);\n"
32 " int tid_y = get_global_id(1);\n"
33 " float4 color;\n"
34 "\n"
35 " color = read_imagef(srcimg, sampler, (int2)(tid_x, tid_y));\n"
36 " write_imagef(dstimg, (int2)(tid_x, tid_y), color);\n"
37 "\n"
38 "}\n";
39
40 static const char *image_to_image_kernel_float_coord_code =
41 "\n"
42 "__kernel void image_to_image_copy(read_only image2d_t srcimg, write_only image2d_t dstimg, sampler_t sampler)\n"
43 "{\n"
44 " int tid_x = get_global_id(0);\n"
45 " int tid_y = get_global_id(1);\n"
46 " float4 color;\n"
47 "\n"
48 " color = read_imagef(srcimg, sampler, (float2)((float)tid_x, (float)tid_y));\n"
49 " write_imagef(dstimg, (int2)(tid_x, tid_y), color);\n"
50 "\n"
51 "}\n";
52
53
54 static const char *image_sum_kernel_integer_coord_code =
55 "\n"
56 "__kernel void image_sum(read_only image2d_t srcimg0, read_only image2d_t srcimg1, write_only image2d_t dstimg, sampler_t sampler)\n"
57 "{\n"
58 " int tid_x = get_global_id(0);\n"
59 " int tid_y = get_global_id(1);\n"
60 " float4 color0;\n"
61 " float4 color1;\n"
62 "\n"
63 " color0 = read_imagef(srcimg0, sampler, (int2)(tid_x, tid_y));\n"
64 " color1 = read_imagef(srcimg1, sampler, (int2)(tid_x, tid_y));\n"
65 " write_imagef(dstimg, (int2)(tid_x, tid_y), color0 + color1);\n"
66 "\n"
67 "}\n";
68
69
70 static const char *image_sum_kernel_float_coord_code =
71 "\n"
72 "__kernel void image_sum(read_only image2d_t srcimg0, read_only image2d_t srcimg1, write_only image2d_t dstimg, sampler_t sampler)\n"
73 "{\n"
74 " int tid_x = get_global_id(0);\n"
75 " int tid_y = get_global_id(1);\n"
76 " float4 color0;\n"
77 " float4 color1;\n"
78 "\n"
79 " color0 = read_imagef(srcimg0, sampler, (float2)((float)tid_x, (float)tid_y));\n"
80 " color1 = read_imagef(srcimg1, sampler, (float2)((float)tid_x, (float)tid_y));\n"
81 " write_imagef(dstimg,(int2)(tid_x, tid_y), color0 + color1);\n"
82 "\n"
83 "}\n";
84
85
86 static unsigned char *
generate_initial_byte_image(int w,int h,int num_elements,unsigned char value)87 generate_initial_byte_image(int w, int h, int num_elements, unsigned char value)
88 {
89 unsigned char *ptr = (unsigned char*)malloc(w * h * num_elements);
90 int i;
91
92 for (i = 0; i < w*h*num_elements; i++)
93 ptr[i] = value;
94
95 return ptr;
96 }
97
98 static unsigned char *
generate_expected_byte_image(unsigned char ** input_data,int num_inputs,int w,int h,int num_elements)99 generate_expected_byte_image(unsigned char **input_data, int num_inputs, int w, int h, int num_elements)
100 {
101 unsigned char *ptr = (unsigned char*)malloc(w * h * num_elements);
102 int i;
103
104 for (i = 0; i < w*h*num_elements; i++)
105 {
106 int j;
107 ptr[i] = 0;
108 for (j = 0; j < num_inputs; j++)
109 {
110 unsigned char *input = *(input_data + j);
111 ptr[i] += input[i];
112 }
113 }
114
115 return ptr;
116 }
117
118
119 static unsigned char *
generate_byte_image(int w,int h,int num_elements,MTdata d)120 generate_byte_image(int w, int h, int num_elements, MTdata d)
121 {
122 unsigned char *ptr = (unsigned char*)malloc(w * h * num_elements);
123 int i;
124
125 for (i = 0; i < w*h*num_elements; i++)
126 ptr[i] = (unsigned char)genrand_int32(d) & 31;
127
128 return ptr;
129 }
130
131 static int
verify_byte_image(unsigned char * image,unsigned char * outptr,int w,int h,int num_elements)132 verify_byte_image(unsigned char *image, unsigned char *outptr, int w, int h, int num_elements)
133 {
134 int i;
135
136 for (i = 0; i < w*h*num_elements; i++)
137 {
138 if (outptr[i] != image[i])
139 {
140 return -1;
141 }
142 }
143 return 0;
144 }
145
146 int
test_image_multipass_integer_coord(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)147 test_image_multipass_integer_coord(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
148 {
149 int img_width = 512;
150 int img_height = 512;
151 cl_image_format img_format;
152
153 int num_input_streams = 8;
154 cl_mem *input_streams;
155 cl_mem accum_streams[2];
156 unsigned char *expected_output;
157 unsigned char *output_ptr;
158 cl_kernel kernel[2];
159 int err;
160
161 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
162
163 img_format.image_channel_order = CL_RGBA;
164 img_format.image_channel_data_type = CL_UNORM_INT8;
165
166 expected_output = (unsigned char*)malloc(sizeof(unsigned char) * 4 * img_width * img_height);
167 output_ptr = (unsigned char*)malloc(sizeof(unsigned char) * 4 * img_width * img_height);
168
169 // Create the accum images with initial data.
170 {
171 unsigned char *initial_data;
172 cl_mem_flags flags;
173
174 initial_data = generate_initial_byte_image(img_width, img_height, 4, 0xF0);
175 flags = CL_MEM_READ_WRITE;
176
177 accum_streams[0] = create_image_2d(context, flags, &img_format, img_width, img_height, 0, NULL, NULL);
178 if (!accum_streams[0])
179 {
180 log_error("create_image_2d failed\n");
181 free(expected_output);
182 free(output_ptr);
183 return -1;
184 }
185
186 size_t origin[3] = {0, 0, 0}, region[3] = {img_width, img_height, 1};
187 err = clEnqueueWriteImage(queue, accum_streams[0], CL_TRUE,
188 origin, region, 0, 0,
189 initial_data, 0, NULL, NULL);
190 if (err)
191 {
192 log_error("clWriteImage failed: %d\n", err);
193 free(expected_output);
194 free(output_ptr);
195 return -1;
196 }
197
198 accum_streams[1] = create_image_2d(context, flags, &img_format, img_width, img_height, 0, NULL, NULL);
199 if (!accum_streams[1])
200 {
201 log_error("create_image_2d failed\n");
202 free(expected_output);
203 free(output_ptr);
204 return -1;
205 }
206 err = clEnqueueWriteImage(queue, accum_streams[1], CL_TRUE,
207 origin, region, 0, 0,
208 initial_data, 0, NULL, NULL);
209 if (err)
210 {
211 log_error("clWriteImage failed: %d\n", err);
212 free(expected_output);
213 free(output_ptr);
214 return -1;
215 }
216
217 free(initial_data);
218 }
219
220 // Set up the input data.
221 {
222 cl_mem_flags flags;
223 unsigned char **input_data = (unsigned char **)malloc(sizeof(unsigned char*) * num_input_streams);
224 MTdata d;
225
226 input_streams = (cl_mem*)malloc(sizeof(cl_mem) * num_input_streams);
227 flags = CL_MEM_READ_WRITE;
228
229 int i;
230 d = init_genrand( gRandomSeed );
231 for ( i = 0; i < num_input_streams; i++)
232 {
233 input_data[i] = generate_byte_image(img_width, img_height, 4, d);
234 input_streams[i] = create_image_2d(context, flags, &img_format, img_width, img_height, 0, NULL, NULL);
235 if (!input_streams[i])
236 {
237 log_error("create_image_2d failed\n");
238 free_mtdata(d);
239 free(expected_output);
240 free(output_ptr);
241 return -1;
242 }
243
244 size_t origin[3] = {0, 0, 0}, region[3] = {img_width, img_height, 1};
245 err = clEnqueueWriteImage(queue, input_streams[i], CL_TRUE,
246 origin, region, 0, 0,
247 input_data[i], 0, NULL, NULL);
248 if (err)
249 {
250 log_error("clWriteImage failed: %d\n", err);
251 free_mtdata(d);
252 free(expected_output);
253 free(output_ptr);
254 free(input_streams);
255 return -1;
256 }
257
258
259 }
260 free_mtdata(d); d = NULL;
261 expected_output = generate_expected_byte_image(input_data, num_input_streams, img_width, img_height, 4);
262 for ( i = 0; i < num_input_streams; i++)
263 {
264 free(input_data[i]);
265 }
266 free( input_data );
267 }
268
269 // Set up the kernels.
270 {
271 cl_program program[4];
272
273 err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &image_to_image_kernel_integer_coord_code, "image_to_image_copy");
274 if (err)
275 {
276 log_error("Failed to create kernel 0: %d\n", err);
277 return -1;
278 }
279 err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &image_sum_kernel_integer_coord_code, "image_sum");
280 if (err)
281 {
282 log_error("Failed to create kernel 1: %d\n", err);
283 return -1;
284 }
285 clReleaseProgram(program[0]);
286 clReleaseProgram(program[1]);
287 }
288
289 cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err);
290 test_error(err, "clCreateSampler failed");
291
292 {
293 size_t threads[3] = {0, 0, 0};
294 threads[0] = (size_t)img_width;
295 threads[1] = (size_t)img_height;
296 int i;
297
298 {
299 cl_mem accum_input;
300 cl_mem accum_output;
301
302 err = clSetKernelArg(kernel[0], 0, sizeof input_streams[0], &input_streams[0]);
303 err |= clSetKernelArg(kernel[0], 1, sizeof accum_streams[0], &accum_streams[0]);
304 err |= clSetKernelArg(kernel[0], 2, sizeof sampler, &sampler);
305 if (err != CL_SUCCESS)
306 {
307 log_error("clSetKernelArgs failed\n");
308 return -1;
309 }
310 err = clEnqueueNDRangeKernel( queue, kernel[0], 2, NULL, threads, NULL, 0, NULL, NULL );
311 if (err != CL_SUCCESS)
312 {
313 log_error("clEnqueueNDRangeKernel failed\n");
314 return -1;
315 }
316
317 for (i = 1; i < num_input_streams; i++)
318 {
319 accum_input = accum_streams[(i-1)%2];
320 accum_output = accum_streams[i%2];
321
322 err = clSetKernelArg(kernel[1], 0, sizeof accum_input, &accum_input);
323 err |= clSetKernelArg(kernel[1], 1, sizeof input_streams[i], &input_streams[i]);
324 err |= clSetKernelArg(kernel[1], 2, sizeof accum_output, &accum_output);
325 err |= clSetKernelArg(kernel[1], 3, sizeof sampler, &sampler);
326
327 if (err != CL_SUCCESS)
328 {
329 log_error("clSetKernelArgs failed\n");
330 return -1;
331 }
332 err = clEnqueueNDRangeKernel( queue, kernel[1], 2, NULL, threads, NULL, 0, NULL, NULL );
333 if (err != CL_SUCCESS)
334 {
335 log_error("clEnqueueNDRangeKernel failed\n");
336 return -1;
337 }
338 }
339
340 // Copy the last accum into the other one.
341 accum_input = accum_streams[(i-1)%2];
342 accum_output = accum_streams[i%2];
343 err = clSetKernelArg(kernel[0], 0, sizeof accum_input, &accum_input);
344 err |= clSetKernelArg(kernel[0], 1, sizeof accum_output, &accum_output);
345 if (err != CL_SUCCESS)
346 {
347 log_error("clSetKernelArgs failed\n");
348 return -1;
349 }
350 err = clEnqueueNDRangeKernel( queue, kernel[0], 2, NULL, threads, NULL, 0, NULL, NULL );
351 if (err != CL_SUCCESS)
352 {
353 log_error("clEnqueueNDRangeKernel failed\n");
354 return -1;
355 }
356
357 size_t origin[3] = {0, 0, 0}, region[3] = {img_width, img_height, 1};
358 err = clEnqueueReadImage(queue, accum_output, CL_TRUE,
359 origin, region, 0, 0,
360 (void *)output_ptr, 0, NULL, NULL);
361 if (err != CL_SUCCESS)
362 {
363 log_error("clReadImage failed\n");
364 return -1;
365 }
366 err = verify_byte_image(expected_output, output_ptr, img_width, img_height, 4);
367 if (err)
368 {
369 log_error("IMAGE_MULTIPASS test failed.\n");
370 }
371 else
372 {
373 log_info("IMAGE_MULTIPASS test passed\n");
374 }
375 }
376
377 clReleaseSampler(sampler);
378 }
379
380
381 // cleanup
382 clReleaseMemObject(accum_streams[0]);
383 clReleaseMemObject(accum_streams[1]);
384 {
385 int i;
386 for (i = 0; i < num_input_streams; i++)
387 {
388 clReleaseMemObject(input_streams[i]);
389 }
390 }
391 free(input_streams);
392 clReleaseKernel(kernel[0]);
393 clReleaseKernel(kernel[1]);
394 free(expected_output);
395 free(output_ptr);
396
397 return err;
398 }
399
400 int
test_image_multipass_float_coord(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)401 test_image_multipass_float_coord(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
402 {
403 int img_width = 512;
404 int img_height = 512;
405 cl_image_format img_format;
406
407 int num_input_streams = 8;
408 cl_mem *input_streams;
409 cl_mem accum_streams[2];
410 unsigned char *expected_output;
411 unsigned char *output_ptr;
412 cl_kernel kernel[2];
413 int err;
414
415 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
416
417 img_format.image_channel_order = CL_RGBA;
418 img_format.image_channel_data_type = CL_UNORM_INT8;
419
420 output_ptr = (unsigned char*)malloc(sizeof(unsigned char) * 4 * img_width * img_height);
421
422 // Create the accum images with initial data.
423 {
424 unsigned char *initial_data;
425 cl_mem_flags flags;
426
427 initial_data = generate_initial_byte_image(img_width, img_height, 4, 0xF0);
428 flags = CL_MEM_READ_WRITE;
429
430 accum_streams[0] = create_image_2d(context, flags, &img_format, img_width, img_height, 0, NULL, NULL);
431 if (!accum_streams[0])
432 {
433 log_error("create_image_2d failed\n");
434 return -1;
435 }
436
437 size_t origin[3] = {0, 0, 0}, region[3] = {img_width, img_height, 1};
438 err = clEnqueueWriteImage(queue, accum_streams[0], CL_TRUE,
439 origin, region, 0, 0,
440 initial_data, 0, NULL, NULL);
441 if (err)
442 {
443 log_error("clWriteImage failed: %d\n", err);
444 return -1;
445 }
446
447 accum_streams[1] = create_image_2d(context, flags, &img_format, img_width, img_height, 0, NULL, NULL);
448 if (!accum_streams[1])
449 {
450 log_error("create_image_2d failed\n");
451 return -1;
452 }
453 err = clEnqueueWriteImage(queue, accum_streams[1], CL_TRUE,
454 origin, region, 0, 0,
455 initial_data, 0, NULL, NULL);
456 if (err)
457 {
458 log_error("clWriteImage failed: %d\n", err);
459 return -1;
460 }
461
462 free(initial_data);
463 }
464
465 // Set up the input data.
466 {
467 cl_mem_flags flags;
468 unsigned char **input_data = (unsigned char **)malloc(sizeof(unsigned char*) * num_input_streams);
469 MTdata d;
470
471 input_streams = (cl_mem*)malloc(sizeof(cl_mem) * num_input_streams);
472 flags = CL_MEM_READ_WRITE;
473
474 int i;
475 d = init_genrand( gRandomSeed );
476 for ( i = 0; i < num_input_streams; i++)
477 {
478 input_data[i] = generate_byte_image(img_width, img_height, 4, d);
479 input_streams[i] = create_image_2d(context, flags, &img_format, img_width, img_height, 0, NULL, NULL);
480 if (!input_streams[i])
481 {
482 log_error("create_image_2d failed\n");
483 free(input_data);
484 free(input_streams);
485 return -1;
486 }
487
488 size_t origin[3] = {0, 0, 0}, region[3] = {img_width, img_height, 1};
489 err = clEnqueueWriteImage(queue, input_streams[i], CL_TRUE,
490 origin, region, 0, 0,
491 input_data[i], 0, NULL, NULL);
492 if (err)
493 {
494 log_error("clWriteImage failed: %d\n", err);
495 free(input_data);
496 free(input_streams);
497 return -1;
498 }
499 }
500 free_mtdata(d); d = NULL;
501 expected_output = generate_expected_byte_image(input_data, num_input_streams, img_width, img_height, 4);
502 for ( i = 0; i < num_input_streams; i++)
503 {
504 free(input_data[i]);
505 }
506 free(input_data);
507 }
508
509 // Set up the kernels.
510 {
511 cl_program program[2];
512
513 err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, &image_to_image_kernel_float_coord_code, "image_to_image_copy");
514 if (err)
515 {
516 log_error("Failed to create kernel 2: %d\n", err);
517 return -1;
518 }
519 err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, &image_sum_kernel_float_coord_code, "image_sum");
520 if (err)
521 {
522 log_error("Failed to create kernel 3: %d\n", err);
523 return -1;
524 }
525
526 clReleaseProgram(program[0]);
527 clReleaseProgram(program[1]);
528 }
529
530 cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err);
531 test_error(err, "clCreateSampler failed");
532
533 {
534 size_t threads[3] = {0, 0, 0};
535 threads[0] = (size_t)img_width;
536 threads[1] = (size_t)img_height;
537 int i;
538
539 {
540 cl_mem accum_input;
541 cl_mem accum_output;
542
543 err = clSetKernelArg(kernel[0], 0, sizeof input_streams[0], &input_streams[0]);
544 err |= clSetKernelArg(kernel[0], 1, sizeof accum_streams[0], &accum_streams[0]);
545 err |= clSetKernelArg(kernel[0], 2, sizeof sampler, &sampler);
546 if (err != CL_SUCCESS)
547 {
548 log_error("clSetKernelArgs failed\n");
549 return -1;
550 }
551 err = clEnqueueNDRangeKernel( queue, kernel[0], 2, NULL, threads, NULL, 0, NULL, NULL );
552 if (err != CL_SUCCESS)
553 {
554 log_error("clEnqueueNDRangeKernel failed\n");
555 return -1;
556 }
557
558 for (i = 1; i < num_input_streams; i++)
559 {
560 accum_input = accum_streams[(i-1)%2];
561 accum_output = accum_streams[i%2];
562
563 err = clSetKernelArg(kernel[1], 0, sizeof accum_input, &accum_input);
564 err |= clSetKernelArg(kernel[1], 1, sizeof input_streams[i], &input_streams[i]);
565 err |= clSetKernelArg(kernel[1], 2, sizeof accum_output, &accum_output);
566 err |= clSetKernelArg(kernel[1], 3, sizeof sampler, &sampler);
567
568 if (err != CL_SUCCESS)
569 {
570 log_error("clSetKernelArgs failed\n");
571 return -1;
572 }
573 err = clEnqueueNDRangeKernel( queue, kernel[1], 2, NULL, threads, NULL, 0, NULL, NULL );
574 if (err != CL_SUCCESS)
575 {
576 log_error("clEnqueueNDRangeKernel failed\n");
577 return -1;
578 }
579 }
580
581 // Copy the last accum into the other one.
582 accum_input = accum_streams[(i-1)%2];
583 accum_output = accum_streams[i%2];
584 err = clSetKernelArg(kernel[0], 0, sizeof accum_input, &accum_input);
585 err |= clSetKernelArg(kernel[0], 1, sizeof accum_output, &accum_output);
586 if (err != CL_SUCCESS)
587 {
588 log_error("clSetKernelArgs failed\n");
589 return -1;
590 }
591 err = clEnqueueNDRangeKernel( queue, kernel[0], 2, NULL, threads, NULL, 0, NULL, NULL );
592 if (err != CL_SUCCESS)
593 {
594 log_error("clEnqueueNDRangeKernel failed\n");
595 return -1;
596 }
597
598 size_t origin[3] = {0, 0, 0}, region[3] = {img_width, img_height, 1};
599 err = clEnqueueReadImage(queue, accum_output, CL_TRUE,
600 origin, region, 0, 0,
601 (void *)output_ptr, 0, NULL, NULL);
602 if (err != CL_SUCCESS)
603 {
604 log_error("clReadImage failed\n");
605 return -1;
606 }
607 err = verify_byte_image(expected_output, output_ptr, img_width, img_height, 4);
608 if (err)
609 {
610 log_error("IMAGE_MULTIPASS test failed.\n");
611 }
612 else
613 {
614 log_info("IMAGE_MULTIPASS test passed\n");
615 }
616 }
617
618 }
619
620
621 // cleanup
622 clReleaseSampler(sampler);
623 clReleaseMemObject(accum_streams[0]);
624 clReleaseMemObject(accum_streams[1]);
625 {
626 int i;
627 for (i = 0; i < num_input_streams; i++)
628 {
629 clReleaseMemObject(input_streams[i]);
630 }
631 }
632 clReleaseKernel(kernel[0]);
633 clReleaseKernel(kernel[1]);
634 free(expected_output);
635 free(output_ptr);
636 free(input_streams);
637
638 return err;
639 }
640
641
642
643
644
645