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 *multireadimage_kernel_code =
28 "__kernel void test_multireadimage(int n, int m, sampler_t sampler, \n"
29 " read_only image2d_t img0, read_only image2d_t img1, \n"
30 " read_only image2d_t img2, read_only image2d_t img3, \n"
31 " read_only image2d_t img4, read_only image2d_t img5, \n"
32 " read_only image2d_t img6, __global float4 *dst)\n"
33 "{\n"
34 " int tid_x = get_global_id(0);\n"
35 " int tid_y = get_global_id(1);\n"
36 " int2 tid = (int2)(tid_x, tid_y);\n"
37 " int indx = tid_y * get_image_width(img5) + tid_x;\n"
38 " float4 sum;\n"
39 "\n"
40 " sum = read_imagef(img0, sampler, tid);\n"
41 " sum += read_imagef(img1, sampler, tid);\n"
42 " sum += read_imagef(img2, sampler, tid);\n"
43 " sum += read_imagef(img3, sampler, tid);\n"
44 " sum += read_imagef(img4, sampler, tid);\n"
45 " sum += read_imagef(img5, sampler, tid);\n"
46 " sum += read_imagef(img6, sampler, tid);\n"
47 "\n"
48 " dst[indx] = sum;\n"
49 "}\n";
50
51
52 static unsigned char *
generate_8888_image(int w,int h,MTdata d)53 generate_8888_image(int w, int h, MTdata d)
54 {
55 unsigned char *ptr = (unsigned char*)malloc(w * h * 4);
56 int i;
57
58 for (i=0; i<w*h*4; i++)
59 ptr[i] = (unsigned char)genrand_int32(d);
60
61 return ptr;
62 }
63
64 static int
verify_multireadimage(void * image[],int num_images,float * outptr,int w,int h)65 verify_multireadimage(void *image[], int num_images, float *outptr, int w, int h)
66 {
67 int i, j;
68 float sum;
69 float ulp, max_ulp = 0.0f;
70
71 // ULP error of 1.5 for each read_imagef plus 0.5 for each addition.
72 float max_ulp_allowed = (float)(num_images*1.5+0.5*(num_images-1));
73
74 for (i=0; i<w*h*4; i++)
75 {
76 sum = 0.0f;
77 for (j=0; j<num_images; j++)
78 {
79 sum += ((float)((unsigned char *)image[j])[i] / 255.0f);
80 }
81 ulp = Ulp_Error(outptr[i], sum);
82 if (ulp > max_ulp)
83 max_ulp = ulp;
84 }
85
86 if (max_ulp > max_ulp_allowed)
87 {
88 log_error("READ_MULTIREADIMAGE_RGBA8888 test failed. Max ULP err = %g\n", max_ulp);
89 return -1;
90 }
91 log_info("READ_MULTIREADIMAGE_RGBA8888 test passed. Max ULP err = %g\n", max_ulp);
92 return 0;
93 }
94
95
test_mri_one(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)96 int test_mri_one(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
97 {
98 cl_mem streams[8];
99 cl_image_format img_format;
100 void *input_ptr[7], *output_ptr;
101 cl_program program;
102 cl_kernel kernel;
103 size_t threads[2];
104 int img_width = 512;
105 int img_height = 512;
106 int i, err;
107 size_t origin[3] = {0, 0, 0};
108 size_t region[3] = {img_width, img_height, 1};
109 size_t length = img_width * img_height * 4 * sizeof(float);
110 MTdata d;
111
112 PASSIVE_REQUIRE_IMAGE_SUPPORT( device )
113
114 output_ptr = malloc(length);
115
116 d = init_genrand( gRandomSeed );
117 for (i=0; i<7; i++) {
118 input_ptr[i] = (void *)generate_8888_image(img_width, img_height, d);
119
120 img_format.image_channel_order = CL_RGBA;
121 img_format.image_channel_data_type = CL_UNORM_INT8;
122 streams[i] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format, img_width, img_height, 0, NULL, NULL);
123 if (!streams[i])
124 {
125 log_error("create_image_2d failed\n");
126 return -1;
127 }
128
129 err = clEnqueueWriteImage(queue, streams[i], CL_TRUE, origin, region, 0, 0, input_ptr[i], 0, NULL, NULL);
130 if (err != CL_SUCCESS)
131 {
132 log_error("clWriteImage failed\n");
133 return -1;
134 }
135 }
136 free_mtdata(d); d = NULL;
137
138
139 streams[7] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, NULL);
140 if (!streams[7])
141 {
142 log_error("clCreateArray failed\n");
143 return -1;
144 }
145
146 err = create_single_kernel_helper(context, &program, &kernel, 1, &multireadimage_kernel_code, "test_multireadimage");
147 if (err)
148 return -1;
149
150 cl_sampler sampler = clCreateSampler(context, CL_FALSE, CL_ADDRESS_CLAMP_TO_EDGE, CL_FILTER_NEAREST, &err);
151 test_error(err, "clCreateSampler failed");
152
153 err = clSetKernelArg(kernel, 0, sizeof i, &i);
154 err |= clSetKernelArg(kernel, 1, sizeof err, &err);
155 err |= clSetKernelArg(kernel, 2, sizeof sampler, &sampler);
156 for (i=0; i<8; i++)
157 err |= clSetKernelArg(kernel, 3+i, sizeof streams[i], &streams[i]);
158
159 if (err != CL_SUCCESS)
160 {
161 log_error("clSetKernelArgs failed\n");
162 return -1;
163 }
164
165 threads[0] = (unsigned int)img_width;
166 threads[1] = (unsigned int)img_height;
167
168 err = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, threads, NULL, 0, NULL, NULL);
169 if (err != CL_SUCCESS)
170 {
171 log_error("clExecuteKernel failed\n");
172 return -1;
173 }
174 err = clEnqueueReadBuffer(queue, streams[7], CL_TRUE, 0, length, output_ptr, 0, NULL, NULL);
175 if (err != CL_SUCCESS)
176 {
177 log_error("clReadArray failed\n");
178 return -1;
179 }
180
181 err = verify_multireadimage(input_ptr, 7, (float *)output_ptr, img_width, img_height);
182
183 // cleanup
184 clReleaseSampler(sampler);
185 for (i=0; i<8; i++)
186 clReleaseMemObject(streams[i]);
187 clReleaseKernel(kernel);
188 clReleaseProgram(program);
189 for (i=0; i<7; i++)
190 free(input_ptr[i]);
191 free(output_ptr);
192
193 return err;
194 }
195
196
197
198
199
200