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