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 #ifndef TEST_CONFORMANCE_CLCPP_IMAGES_TEST_WRITE_HPP
17 #define TEST_CONFORMANCE_CLCPP_IMAGES_TEST_WRITE_HPP
18
19 #include <algorithm>
20 #include <sstream>
21 #include <string>
22 #include <tuple>
23 #include <vector>
24
25 #include "common.hpp"
26
27
28 namespace test_images_write {
29
30 template<cl_mem_object_type ImageType, cl_channel_type ChannelType>
31 struct image_test : image_test_base<ImageType, ChannelType>
32 {
33 cl_channel_order channel_order;
34
image_testtest_images_write::image_test35 image_test(cl_channel_order channel_order) :
36 channel_order(channel_order)
37 { }
38 // -----------------------------------------------------------------------------------
39 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
40 // -----------------------------------------------------------------------------------
41 #if defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
generate_sourcetest_images_write::image_test42 std::string generate_source()
43 {
44 std::stringstream s;
45 s << R"(
46 typedef )" << type_name<typename image_test::element_type>() << R"( element_type;
47
48 kernel void test(
49 write_only )" << image_test::image_type_name() << R"(_t img,
50 const global int4 *coords,
51 const global element_type *input
52 ) {
53 const ulong gid = get_global_linear_id();
54
55 write_image)" << image_test::function_suffix() <<
56 "(img, coords[gid]." << image_test::coord_accessor() << R"(, input[gid]);
57 }
58 )";
59
60 return s.str();
61 }
62 #else
generate_sourcetest_images_write::image_test63 std::string generate_source()
64 {
65 std::stringstream s;
66 s << R"(
67 #include <opencl_memory>
68 #include <opencl_common>
69 #include <opencl_work_item>
70 #include <opencl_image>
71 using namespace cl;
72 )";
73
74 s << R"(
75 typedef )" << type_name<typename image_test::element_type>() << R"( element_type;
76
77 kernel void test(
78 )" << image_test::image_type_name() << R"(<element_type, image_access::write> img,
79 const global_ptr<int4[]> coords,
80 const global_ptr<element_type[]> input
81 ) {
82 const ulong gid = get_global_linear_id();
83
84 img.write(coords[gid].)" << image_test::coord_accessor() << R"(, input[gid]);
85 }
86 )";
87
88 return s.str();
89 }
90 #endif
91
runtest_images_write::image_test92 int run(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
93 {
94 int error = CL_SUCCESS;
95
96 cl_program program;
97 cl_kernel kernel;
98
99 std::string kernel_name = "test";
100 std::string source = generate_source();
101
102 // -----------------------------------------------------------------------------------
103 // ------------- ONLY FOR OPENCL 22 CONFORMANCE TEST 22 DEVELOPMENT ------------------
104 // -----------------------------------------------------------------------------------
105 // Only OpenCL C++ to SPIR-V compilation
106 #if defined(DEVELOPMENT) && defined(ONLY_SPIRV_COMPILATION)
107 error = create_opencl_kernel(
108 context, &program, &kernel,
109 source, kernel_name
110 );
111 RETURN_ON_ERROR(error)
112 return error;
113 // Use OpenCL C kernels instead of OpenCL C++ kernels (test C++ host code)
114 #elif defined(DEVELOPMENT) && defined(USE_OPENCLC_KERNELS)
115 error = create_opencl_kernel(
116 context, &program, &kernel,
117 source, kernel_name, "-cl-std=CL2.0", false
118 );
119 RETURN_ON_ERROR(error)
120 // Normal run
121 #else
122 error = create_opencl_kernel(
123 context, &program, &kernel,
124 source, kernel_name
125 );
126 RETURN_ON_ERROR(error)
127 #endif
128
129 using element_type = typename image_test::element_type;
130 using coord_type = cl_int4;
131 using scalar_element_type = typename scalar_type<element_type>::type;
132 using channel_type = typename image_test::channel_type;
133
134 cl_image_format image_format;
135 image_format.image_channel_order = channel_order;
136 image_format.image_channel_data_type = ChannelType;
137
138 const size_t pixel_size = get_pixel_size(&image_format);
139 const size_t channel_count = get_channel_order_channel_count(image_format.image_channel_order);
140
141 cl_image_desc image_desc;
142 image_desc.image_type = ImageType;
143 if (ImageType == CL_MEM_OBJECT_IMAGE1D)
144 {
145 image_desc.image_width = 2048;
146 image_desc.image_height = 1;
147 image_desc.image_depth = 1;
148 }
149 else if (ImageType == CL_MEM_OBJECT_IMAGE2D)
150 {
151 image_desc.image_width = 256;
152 image_desc.image_height = 256;
153 image_desc.image_depth = 1;
154 }
155 else if (ImageType == CL_MEM_OBJECT_IMAGE3D)
156 {
157 image_desc.image_width = 64;
158 image_desc.image_height = 64;
159 image_desc.image_depth = 64;
160 }
161 image_desc.image_array_size = 0;
162 image_desc.image_row_pitch = image_desc.image_width * pixel_size;
163 image_desc.image_slice_pitch = image_desc.image_row_pitch * image_desc.image_height;
164 image_desc.num_mip_levels = 0;
165 image_desc.num_samples = 0;
166 image_desc.mem_object = NULL;
167
168 image_descriptor image_info = create_image_descriptor(image_desc, &image_format);
169
170 std::vector<channel_type> random_image_values = generate_input(
171 image_desc.image_width * image_desc.image_height * image_desc.image_depth * channel_count,
172 image_test::channel_min(), image_test::channel_max(),
173 std::vector<channel_type>()
174 );
175
176 const size_t count = num_elements;
177
178 std::vector<coord_type> coords = generate_input(
179 count,
180 detail::make_value<coord_type>(0),
181 coord_type {
182 static_cast<cl_int>(image_desc.image_width - 1),
183 static_cast<cl_int>(image_desc.image_height - 1),
184 static_cast<cl_int>(image_desc.image_depth - 1),
185 0
186 },
187 std::vector<coord_type>()
188 );
189
190 std::vector<element_type> input(count);
191 for (size_t i = 0; i < count; i++)
192 {
193 const coord_type c = coords[i];
194
195 // Use read_image_pixel from harness/imageHelpers to fill input values
196 // (it will deal with correct channels, orders etc.)
197 read_image_pixel<scalar_element_type>(static_cast<void *>(random_image_values.data()), &image_info,
198 c.s[0], c.s[1], c.s[2],
199 input[i].s);
200 }
201
202 // image_row_pitch and image_slice_pitch must be 0, when clCreateImage is used with host_ptr = NULL
203 image_desc.image_row_pitch = 0;
204 image_desc.image_slice_pitch = 0;
205 cl_mem img = clCreateImage(context, CL_MEM_WRITE_ONLY,
206 &image_format, &image_desc, NULL, &error);
207 RETURN_ON_CL_ERROR(error, "clCreateImage")
208
209 cl_mem coords_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
210 sizeof(coord_type) * count, static_cast<void *>(coords.data()), &error);
211 RETURN_ON_CL_ERROR(error, "clCreateBuffer")
212
213 cl_mem input_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR,
214 sizeof(element_type) * count, static_cast<void *>(input.data()), &error);
215 RETURN_ON_CL_ERROR(error, "clCreateBuffer")
216
217 error = clSetKernelArg(kernel, 0, sizeof(cl_mem), &img);
218 RETURN_ON_CL_ERROR(error, "clSetKernelArg")
219 error = clSetKernelArg(kernel, 1, sizeof(coords_buffer), &coords_buffer);
220 RETURN_ON_CL_ERROR(error, "clSetKernelArg")
221 error = clSetKernelArg(kernel, 2, sizeof(input_buffer), &input_buffer);
222 RETURN_ON_CL_ERROR(error, "clSetKernelArg")
223
224 const size_t global_size = count;
225 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, NULL, 0, NULL, NULL);
226 RETURN_ON_CL_ERROR(error, "clEnqueueNDRangeKernel")
227
228 std::vector<channel_type> image_values(image_desc.image_width * image_desc.image_height * image_desc.image_depth * channel_count);
229
230 const size_t origin[3] = { 0 };
231 const size_t region[3] = { image_desc.image_width, image_desc.image_height, image_desc.image_depth };
232 error = clEnqueueReadImage(
233 queue, img, CL_TRUE,
234 origin, region, 0, 0,
235 static_cast<void *>(image_values.data()),
236 0, NULL, NULL
237 );
238 RETURN_ON_CL_ERROR(error, "clEnqueueReadBuffer")
239
240 for (size_t i = 0; i < count; i++)
241 {
242 const coord_type c = coords[i];
243 const element_type expected = input[i];
244
245 element_type result;
246 read_image_pixel<scalar_element_type>(static_cast<void *>(image_values.data()), &image_info,
247 c.s[0], c.s[1], c.s[2],
248 result.s);
249
250 if (!are_equal(result, expected))
251 {
252 RETURN_ON_ERROR_MSG(-1,
253 "Writing to coordinates %s failed. Expected: %s, got: %s",
254 format_value(c).c_str(), format_value(expected).c_str(), format_value(result).c_str()
255 );
256 }
257 }
258
259 clReleaseMemObject(img);
260 clReleaseMemObject(coords_buffer);
261 clReleaseMemObject(input_buffer);
262 clReleaseKernel(kernel);
263 clReleaseProgram(program);
264 return error;
265 }
266 };
267
268 template<cl_mem_object_type ImageType>
run_test_cases(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)269 int run_test_cases(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
270 {
271 if (!is_test_supported(device))
272 return CL_SUCCESS;
273
274 int error = CL_SUCCESS;
275
276 for (auto channel_order : get_channel_orders(device))
277 {
278 error = image_test<ImageType, CL_SIGNED_INT8>(channel_order)
279 .run(device, context, queue, num_elements);
280 RETURN_ON_ERROR(error)
281 error = image_test<ImageType, CL_SIGNED_INT16>(channel_order)
282 .run(device, context, queue, num_elements);
283 RETURN_ON_ERROR(error)
284 error = image_test<ImageType, CL_SIGNED_INT32>(channel_order)
285 .run(device, context, queue, num_elements);
286 RETURN_ON_ERROR(error)
287
288 error = image_test<ImageType, CL_UNSIGNED_INT8>(channel_order)
289 .run(device, context, queue, num_elements);
290 RETURN_ON_ERROR(error)
291 error = image_test<ImageType, CL_UNSIGNED_INT16>(channel_order)
292 .run(device, context, queue, num_elements);
293 RETURN_ON_ERROR(error)
294 error = image_test<ImageType, CL_UNSIGNED_INT32>(channel_order)
295 .run(device, context, queue, num_elements);
296 RETURN_ON_ERROR(error)
297
298 error = image_test<ImageType, CL_FLOAT>(channel_order)
299 .run(device, context, queue, num_elements);
300 RETURN_ON_ERROR(error)
301 }
302
303 return error;
304 }
305
306
AUTO_TEST_CASE(test_images_write_1d)307 AUTO_TEST_CASE(test_images_write_1d)
308 (cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
309 {
310 return run_test_cases<CL_MEM_OBJECT_IMAGE1D>(device, context, queue, num_elements);
311 }
312
AUTO_TEST_CASE(test_images_write_2d)313 AUTO_TEST_CASE(test_images_write_2d)
314 (cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
315 {
316 return run_test_cases<CL_MEM_OBJECT_IMAGE2D>(device, context, queue, num_elements);
317 }
318
AUTO_TEST_CASE(test_images_write_3d)319 AUTO_TEST_CASE(test_images_write_3d)
320 (cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
321 {
322 return run_test_cases<CL_MEM_OBJECT_IMAGE3D>(device, context, queue, num_elements);
323 }
324
325 } // namespace
326
327 #endif // TEST_CONFORMANCE_CLCPP_IMAGES_TEST_WRITE_HPP
328