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 #include "harness/imageHelpers.h"
18
19 #include <stdio.h>
20 #include <stdlib.h>
21 #include <string.h>
22 #include <sys/types.h>
23 #include <sys/stat.h>
24
25
26 #include "procs.h"
27
28 #include <algorithm>
29 #include <string>
30 #include <vector>
31
32 #include "procs.h"
33
34 namespace {
35 const char *kernel_source = R"(
36 __kernel void test_CL_BGRACL_UNORM_INT8(__global unsigned char *src, write_only image2d_t dstimg)
37 {
38 int tid_x = get_global_id(0);
39 int tid_y = get_global_id(1);
40 int indx = tid_y * get_image_width(dstimg) + tid_x;
41 float4 color;
42
43 indx *= 4;
44 color = (float4)((float)src[indx+2], (float)src[indx+1], (float)src[indx+0], (float)src[indx+3]);
45 color /= (float4)(255.0f, 255.0f, 255.0f, 255.0f);
46 write_imagef(dstimg, (int2)(tid_x, tid_y), color);
47 }
48
49 __kernel void test_CL_RGBACL_UNORM_INT8(__global unsigned char *src, write_only image2d_t dstimg)
50 {
51 int tid_x = get_global_id(0);
52 int tid_y = get_global_id(1);
53 int indx = tid_y * get_image_width(dstimg) + tid_x;
54 float4 color;
55
56 indx *= 4;
57 color = (float4)((float)src[indx+0], (float)src[indx+1], (float)src[indx+2], (float)src[indx+3]);
58 color /= (float4)(255.0f, 255.0f, 255.0f, 255.0f);
59 write_imagef(dstimg, (int2)(tid_x, tid_y), color);
60 }
61
62 __kernel void test_CL_RGBACL_UNORM_INT16(__global unsigned short *src, write_only image2d_t dstimg)
63 {
64 int tid_x = get_global_id(0);
65 int tid_y = get_global_id(1);
66 int indx = tid_y * get_image_width(dstimg) + tid_x;
67 float4 color;
68
69 indx *= 4;
70 color = (float4)((float)src[indx+0], (float)src[indx+1], (float)src[indx+2], (float)src[indx+3]);
71 color /= 65535.0f;
72 write_imagef(dstimg, (int2)(tid_x, tid_y), color);
73 }
74
75 __kernel void test_CL_RGBACL_FLOAT(__global float *src, write_only image2d_t dstimg)
76 {
77 int tid_x = get_global_id(0);
78 int tid_y = get_global_id(1);
79 int indx = tid_y * get_image_width(dstimg) + tid_x;
80 float4 color;
81
82 indx *= 4;
83 color = (float4)(src[indx+0], src[indx+1], src[indx+2], src[indx+3]);
84 write_imagef(dstimg, (int2)(tid_x, tid_y), color);
85 }
86 )";
87
88
generate_random_inputs(std::vector<T> & v)89 template <typename T> void generate_random_inputs(std::vector<T> &v)
90 {
91 RandomSeed seed(gRandomSeed);
92
93 auto random_generator = [&seed]() {
94 return static_cast<T>(genrand_int32(seed));
95 };
96
97 std::generate(v.begin(), v.end(), random_generator);
98 }
99
generate_random_inputs(std::vector<float> & v)100 template <> void generate_random_inputs<float>(std::vector<float> &v)
101 {
102 RandomSeed seed(gRandomSeed);
103
104 auto random_generator = [&seed]() {
105 return get_random_float(-0x40000000, 0x40000000, seed);
106 };
107
108 std::generate(v.begin(), v.end(), random_generator);
109 }
110
111
get_mem_flag_name(cl_mem_flags flags)112 const char *get_mem_flag_name(cl_mem_flags flags)
113 {
114 switch (flags)
115 {
116 case CL_MEM_READ_WRITE: return "CL_MEM_READ_WRITE";
117 case CL_MEM_WRITE_ONLY: return "CL_MEM_WRITE_ONLY";
118 default: return "Unsupported cl_mem_flags value";
119 }
120 }
121
122 template <typename T>
test_writeimage(cl_device_id device,cl_context context,cl_command_queue queue,const cl_image_format * img_format,cl_mem_flags img_flags)123 int test_writeimage(cl_device_id device, cl_context context,
124 cl_command_queue queue, const cl_image_format *img_format,
125 cl_mem_flags img_flags)
126 {
127 clMemWrapper streams[2];
128 clProgramWrapper program;
129 clKernelWrapper kernel;
130
131 std::string kernel_name("test_");
132
133 size_t img_width = 512;
134 size_t img_height = 512;
135
136 int err;
137
138 const size_t origin[3] = { 0, 0, 0 };
139 const size_t region[3] = { img_width, img_height, 1 };
140
141 const size_t num_elements = img_width * img_height * 4;
142 const size_t length = num_elements * sizeof(T);
143
144 PASSIVE_REQUIRE_IMAGE_SUPPORT(device)
145
146 std::vector<T> input(num_elements);
147 std::vector<T> output(num_elements);
148
149 generate_random_inputs(input);
150
151 streams[0] = create_image_2d(context, img_flags, img_format, img_width,
152 img_height, 0, nullptr, &err);
153 test_error(err, "create_image failed.");
154
155 streams[1] =
156 clCreateBuffer(context, CL_MEM_READ_WRITE, length, nullptr, &err);
157 test_error(err, "clCreateBuffer failed.");
158
159 err = clEnqueueWriteBuffer(queue, streams[1], CL_TRUE, 0, length,
160 input.data(), 0, nullptr, nullptr);
161 test_error(err, "clEnqueueWriteImage failed.");
162
163 kernel_name += GetChannelOrderName(img_format->image_channel_order);
164 kernel_name += GetChannelTypeName(img_format->image_channel_data_type);
165
166 err = create_single_kernel_helper(context, &program, &kernel, 1,
167 &kernel_source, kernel_name.c_str());
168 test_error(err, "create_single_kernel_helper failed.");
169
170 err |= clSetKernelArg(kernel, 0, sizeof(streams[1]), &streams[1]);
171 err |= clSetKernelArg(kernel, 1, sizeof(streams[0]), &streams[0]);
172 test_error(err, "clSetKernelArgs failed\n");
173
174 size_t threads[] = { img_width, img_height };
175 err = clEnqueueNDRangeKernel(queue, kernel, 2, nullptr, threads, nullptr, 0,
176 nullptr, nullptr);
177 test_error(err, "clEnqueueNDRangeKernel failed\n");
178
179 err = clEnqueueReadImage(queue, streams[0], CL_TRUE, origin, region, 0, 0,
180 output.data(), 0, nullptr, nullptr);
181
182 if (0 != memcmp(input.data(), output.data(), length))
183 {
184 log_error("WRITE_IMAGE_%s_%s with %s test failed\n",
185 GetChannelOrderName(img_format->image_channel_order),
186 GetChannelTypeName(img_format->image_channel_data_type),
187 get_mem_flag_name(img_flags));
188 err = -1;
189 }
190 else
191 {
192 log_info("WRITE_IMAGE_%s_%s with %s test passed\n",
193 GetChannelOrderName(img_format->image_channel_order),
194 GetChannelTypeName(img_format->image_channel_data_type),
195 get_mem_flag_name(img_flags));
196 }
197
198 return err;
199 }
200
check_format(cl_device_id device,cl_context context,cl_mem_object_type image_type,const cl_image_format img_format,cl_mem_flags test_flags)201 bool check_format(cl_device_id device, cl_context context,
202 cl_mem_object_type image_type,
203 const cl_image_format img_format, cl_mem_flags test_flags)
204 {
205 return is_image_format_required(img_format, test_flags, image_type, device)
206 || is_image_format_supported(context, test_flags, image_type,
207 &img_format);
208 }
209 }
test_writeimage(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)210 int test_writeimage(cl_device_id device, cl_context context,
211 cl_command_queue queue, int num_elements)
212 {
213 int err = 0;
214 const cl_image_format format[] = { { CL_RGBA, CL_UNORM_INT8 },
215 { CL_BGRA, CL_UNORM_INT8 } };
216 const cl_mem_flags test_flags[] = { CL_MEM_WRITE_ONLY, CL_MEM_READ_WRITE };
217
218 for (size_t i = 0; i < ARRAY_SIZE(test_flags) && !err; i++)
219 {
220 err = test_writeimage<cl_uchar>(device, context, queue, &format[0],
221 test_flags[i]);
222
223 if (check_format(device, context, CL_MEM_OBJECT_IMAGE2D, format[1],
224 test_flags[i]))
225 {
226 err |= test_writeimage<cl_uchar>(device, context, queue, &format[1],
227 test_flags[i]);
228 }
229 }
230 return err;
231 }
232
test_writeimage_int16(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)233 int test_writeimage_int16(cl_device_id device, cl_context context,
234 cl_command_queue queue, int num_elements)
235 {
236 int err = 0;
237 const cl_image_format format = { CL_RGBA, CL_UNORM_INT16 };
238 const cl_mem_flags test_flags[] = { CL_MEM_WRITE_ONLY, CL_MEM_READ_WRITE };
239
240 for (size_t i = 0; i < ARRAY_SIZE(test_flags) && !err; i++)
241 {
242 err = test_writeimage<cl_ushort>(device, context, queue, &format,
243 test_flags[i]);
244 }
245 return err;
246 }
247
test_writeimage_fp32(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)248 int test_writeimage_fp32(cl_device_id device, cl_context context,
249 cl_command_queue queue, int num_elements)
250 {
251 int err = 0;
252 const cl_image_format format = { CL_RGBA, CL_FLOAT };
253 const cl_mem_flags test_flags[] = { CL_MEM_WRITE_ONLY, CL_MEM_READ_WRITE };
254
255 for (size_t i = 0; i < ARRAY_SIZE(test_flags) && !err; i++)
256 {
257 err = test_writeimage<cl_float>(device, context, queue, &format,
258 test_flags[i]);
259 }
260 return err;
261 }
262