1 //
2 // Copyright (c) 2020 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 "testBase.h"
17 #include "harness/propertyHelpers.h"
18 #include "harness/typeWrappers.h"
19 #include <vector>
20 #include <algorithm>
21
22 typedef enum
23 {
24 image,
25 image_with_properties,
26 buffer,
27 buffer_with_properties,
28 subbuffer,
29 } test_type;
30
31 struct test_data
32 {
33 test_type type;
34 std::vector<cl_mem_properties> properties;
35 std::string description;
36 cl_kernel kernel;
37 };
38
create_object_and_check_properties(cl_context context,clMemWrapper & test_object,test_data test_case,cl_mem_flags flags,std::vector<cl_uint> local_data,cl_uint size_x,cl_uint size_y)39 static int create_object_and_check_properties(cl_context context,
40 clMemWrapper& test_object,
41 test_data test_case,
42 cl_mem_flags flags,
43 std::vector<cl_uint> local_data,
44 cl_uint size_x, cl_uint size_y)
45 {
46 cl_int error = CL_SUCCESS;
47
48 switch (test_case.type)
49 {
50 case image: {
51 cl_image_format format = { 0 };
52 format.image_channel_order = CL_RGBA;
53 format.image_channel_data_type = CL_UNSIGNED_INT32;
54 test_object = clCreateImage2D(context, flags, &format, size_x,
55 size_y, 0, local_data.data(), &error);
56 test_error(error, "clCreateImage2D failed");
57 }
58 break;
59 case image_with_properties: {
60 cl_image_format format = { 0 };
61 format.image_channel_order = CL_RGBA;
62 format.image_channel_data_type = CL_UNSIGNED_INT32;
63 cl_image_desc desc = { 0 };
64 desc.image_type = CL_MEM_OBJECT_IMAGE2D;
65 desc.image_width = size_x;
66 desc.image_height = size_y;
67
68 if (test_case.properties.size() == 0)
69 {
70 test_object = clCreateImageWithProperties(
71 context, NULL, flags, &format, &desc, local_data.data(),
72 &error);
73 }
74 else
75 {
76 test_object = clCreateImageWithProperties(
77 context, test_case.properties.data(), flags, &format, &desc,
78 local_data.data(), &error);
79 }
80 test_error(error, "clCreateImageWithProperties failed");
81 }
82 break;
83 case buffer: {
84 test_object = clCreateBuffer(context, flags,
85 local_data.size() * sizeof(cl_uint),
86 local_data.data(), &error);
87 test_error(error, "clCreateBuffer failed");
88 }
89 case buffer_with_properties: {
90 if (test_case.properties.size() == 0)
91 {
92 test_object = clCreateBufferWithProperties(
93 context, NULL, flags, local_data.size() * sizeof(cl_uint),
94 local_data.data(), &error);
95 }
96 else
97 {
98 test_object = clCreateBufferWithProperties(
99 context, test_case.properties.data(), flags,
100 local_data.size() * sizeof(cl_uint), local_data.data(),
101 &error);
102 }
103 test_error(error, "clCreateBufferWithProperties failed.");
104 }
105 break;
106 case subbuffer: {
107 clMemWrapper parent_object;
108 if (test_case.properties.size() == 0)
109 {
110 parent_object = clCreateBufferWithProperties(
111 context, NULL, flags, local_data.size() * sizeof(cl_uint),
112 local_data.data(), &error);
113 }
114 else
115 {
116 parent_object = clCreateBufferWithProperties(
117 context, test_case.properties.data(), flags,
118 local_data.size() * sizeof(cl_uint), local_data.data(),
119 &error);
120 }
121 test_error(error, "clCreateBufferWithProperties failed.");
122
123 cl_mem_flags subbuffer_flags = flags
124 & (CL_MEM_READ_WRITE | CL_MEM_READ_ONLY | CL_MEM_WRITE_ONLY);
125
126 cl_buffer_region region = { 0 };
127 region.origin = 0;
128 region.size = local_data.size() * sizeof(cl_uint);
129 test_object = clCreateSubBuffer(parent_object, subbuffer_flags,
130 CL_BUFFER_CREATE_TYPE_REGION,
131 ®ion, &error);
132 test_error(error, "clCreateSubBuffer failed.");
133 }
134 break;
135 default: log_error("Unknown test type!"); return TEST_FAIL;
136 }
137
138 std::vector<cl_mem_properties> check_properties;
139 size_t set_size = 0;
140
141 error =
142 clGetMemObjectInfo(test_object, CL_MEM_PROPERTIES, 0, NULL, &set_size);
143 test_error(error,
144 "clGetMemObjectInfo failed asking for CL_MEM_PROPERTIES size.");
145
146 // Buffers, subbuffers, and images must return no properties.
147 if (test_case.type == buffer || test_case.type == subbuffer
148 || test_case.type == image)
149 {
150 if (set_size == 0)
151 {
152 return TEST_PASS;
153 }
154 else
155 {
156 log_error("Queried properties must have size equal to zero for "
157 "buffers, subbuffers, and images.");
158 return TEST_FAIL;
159 }
160 }
161
162 if (set_size == 0 && test_case.properties.size() == 0)
163 {
164 return TEST_PASS;
165 }
166 if (set_size != test_case.properties.size() * sizeof(cl_mem_properties))
167 {
168 log_error("ERROR: CL_MEM_PROPERTIES size is %d, expected %d.\n",
169 set_size,
170 test_case.properties.size() * sizeof(cl_queue_properties));
171 return TEST_FAIL;
172 }
173
174 cl_uint number_of_props = set_size / sizeof(cl_mem_properties);
175 check_properties.resize(number_of_props);
176 error = clGetMemObjectInfo(test_object, CL_MEM_PROPERTIES, set_size,
177 check_properties.data(), NULL);
178 test_error(error,
179 "clGetMemObjectInfo failed asking for CL_MEM_PROPERTIES.");
180
181 error = compareProperties(check_properties, test_case.properties);
182 return error;
183 }
184
run_test_query_properties(cl_context context,cl_command_queue queue,test_data test_case)185 static int run_test_query_properties(cl_context context, cl_command_queue queue,
186 test_data test_case)
187 {
188 int error = CL_SUCCESS;
189 log_info("\nTC description: %s\n", test_case.description.c_str());
190
191 clMemWrapper obj_src;
192 clMemWrapper obj_dst;
193 clEventWrapper event;
194 MTdata init_generator = init_genrand(gRandomSeed);
195 cl_mem_flags flags;
196 cl_uint size_x = 4;
197 cl_uint size_y = 4;
198 size_t size = size_x * size_y * 4;
199 size_t global_dim[2] = { size_x, size_y };
200 const size_t origin[3] = { 0, 0, 0 };
201 const size_t region[3] = { size_x, size_y, 1 };
202
203 std::vector<cl_uint> src_data(size);
204 std::vector<cl_uint> dst_data(size);
205
206 generate_random_data(kUInt, size, init_generator, src_data.data());
207 generate_random_data(kUInt, size, init_generator, dst_data.data());
208 free_mtdata(init_generator);
209 init_generator = NULL;
210
211 flags = CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR;
212 error = create_object_and_check_properties(context, obj_src, test_case,
213 flags, src_data, size_x, size_y);
214 test_error(error, "create_object_and_check_properties obj_src failed.");
215
216 flags = CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR;
217 error = create_object_and_check_properties(context, obj_dst, test_case,
218 flags, dst_data, size_x, size_y);
219 test_error(error, "create_object_and_check_properties obj_dst failed.");
220
221 error = clSetKernelArg(test_case.kernel, 0, sizeof(obj_src), &obj_src);
222 test_error(error, "clSetKernelArg 0 failed.");
223
224 error = clSetKernelArg(test_case.kernel, 1, sizeof(obj_dst), &obj_dst);
225 test_error(error, "clSetKernelArg 1 failed.");
226 switch (test_case.type)
227 {
228 case image:
229 case image_with_properties: {
230 error = clEnqueueNDRangeKernel(queue, test_case.kernel, 2, NULL,
231 global_dim, NULL, 0, NULL, &event);
232 test_error(error, "clEnqueueNDRangeKernel failed.");
233
234 error = clWaitForEvents(1, &event);
235 test_error(error, "clWaitForEvents failed.");
236
237 error = clEnqueueReadImage(queue, obj_dst, CL_TRUE, origin, region,
238 0, 0, dst_data.data(), 0, NULL, NULL);
239 test_error(error, "clEnqueueReadImage failed.");
240 }
241 break;
242 case buffer:
243 case buffer_with_properties:
244 case subbuffer: {
245 error = clEnqueueNDRangeKernel(queue, test_case.kernel, 1, NULL,
246 &size, NULL, 0, NULL, &event);
247 test_error(error, "clEnqueueNDRangeKernel failed.");
248
249 error = clWaitForEvents(1, &event);
250 test_error(error, "clWaitForEvents failed.");
251
252 error = clEnqueueReadBuffer(queue, obj_dst, CL_TRUE, 0,
253 dst_data.size() * sizeof(cl_uint),
254 dst_data.data(), 0, NULL, NULL);
255 test_error(error, "clEnqueueReadBuffer failed.");
256 }
257 break;
258 default: log_error("Unknown test type!"); return TEST_FAIL;
259 }
260
261 for (size_t i = 0; i < size; ++i)
262 {
263 if (dst_data[i] != src_data[i])
264 {
265 log_error("ERROR: Output results mismatch.");
266 return TEST_FAIL;
267 }
268 }
269
270 log_info("TC result: passed\n");
271 return TEST_PASS;
272 }
273
test_image_properties_queries(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)274 int test_image_properties_queries(cl_device_id deviceID, cl_context context,
275 cl_command_queue queue, int num_elements)
276 {
277 int error = CL_SUCCESS;
278 cl_bool supports_images = CL_TRUE;
279
280 error = clGetDeviceInfo(deviceID, CL_DEVICE_IMAGE_SUPPORT,
281 sizeof(supports_images), &supports_images, NULL);
282 test_error(error, "clGetDeviceInfo for CL_DEVICE_IMAGE_SUPPORT failed");
283
284 if (supports_images == CL_FALSE)
285 {
286 log_info("No image support on current device - skipped\n");
287 return TEST_SKIPPED_ITSELF;
288 }
289
290 clProgramWrapper program;
291 clKernelWrapper kernel;
292
293 const char* kernel_src = R"CLC(
294 __kernel void data_copy(read_only image2d_t src, write_only image2d_t dst)
295 {
296 int tid_x = get_global_id(0);
297 int tid_y = get_global_id(1);
298 int2 coords = (int2)(tid_x, tid_y);
299 uint4 val = read_imageui(src, coords);
300 write_imageui(dst, coords, val);
301
302 }
303 )CLC";
304
305 error = create_single_kernel_helper(context, &program, &kernel, 1,
306 &kernel_src, "data_copy");
307 test_error(error, "create_single_kernel_helper failed");
308
309 std::vector<test_data> test_cases;
310 test_cases.push_back({ image, {}, "regular image", kernel });
311 test_cases.push_back(
312 { image_with_properties, { 0 }, "image, 0 properties", kernel });
313 test_cases.push_back(
314 { image_with_properties, {}, "image, NULL properties", kernel });
315
316 for (auto test_case : test_cases)
317 {
318 error |= run_test_query_properties(context, queue, test_case);
319 }
320
321 return error;
322 }
323
test_buffer_properties_queries(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)324 int test_buffer_properties_queries(cl_device_id deviceID, cl_context context,
325 cl_command_queue queue, int num_elements)
326 {
327 int error = CL_SUCCESS;
328
329 clProgramWrapper program;
330 clKernelWrapper kernel;
331
332 const char* kernel_src = R"CLC(
333 __kernel void data_copy(__global int *src, __global int *dst)
334 {
335 int tid = get_global_id(0);
336
337 dst[tid] = src[tid];
338
339 }
340 )CLC";
341 error = create_single_kernel_helper(context, &program, &kernel, 1,
342 &kernel_src, "data_copy");
343 test_error(error, "create_single_kernel_helper failed");
344
345 std::vector<test_data> test_cases;
346 test_cases.push_back({ buffer, {}, "regular buffer", kernel });
347 test_cases.push_back(
348 { buffer_with_properties, { 0 }, "buffer with 0 properties", kernel });
349 test_cases.push_back(
350 { buffer_with_properties, {}, "buffer with NULL properties", kernel });
351 test_cases.push_back(
352 { subbuffer, { 0 }, "subbuffer with 0 properties", kernel });
353 test_cases.push_back(
354 { subbuffer, {}, "subbuffer with NULL properties", kernel });
355
356 for (auto test_case : test_cases)
357 {
358 error |= run_test_query_properties(context, queue, test_case);
359 }
360
361 return error;
362 }
363