1 //
2 // Copyright (c) 2022 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 "basic_command_buffer.h"
17 #include "procs.h"
18
19 #include <algorithm>
20 #include <cstring>
21 #include <vector>
22
23
24 //--------------------------------------------------------------------------
BasicCommandBufferTest(cl_device_id device,cl_context context,cl_command_queue queue)25 BasicCommandBufferTest::BasicCommandBufferTest(cl_device_id device,
26 cl_context context,
27 cl_command_queue queue)
28 : CommandBufferTestBase(device), context(context), queue(nullptr),
29 num_elements(0), simultaneous_use_support(false),
30 out_of_order_support(false),
31 // try to use simultaneous path by default
32 simultaneous_use_requested(true),
33 // due to simultaneous cases extend buffer size
34 buffer_size_multiplier(1), command_buffer(this)
35 {
36 cl_int error = clRetainCommandQueue(queue);
37 if (error != CL_SUCCESS)
38 {
39 throw std::runtime_error("clRetainCommandQueue failed\n");
40 }
41 this->queue = queue;
42 }
43
44 //--------------------------------------------------------------------------
Skip()45 bool BasicCommandBufferTest::Skip()
46 {
47 cl_command_queue_properties required_properties;
48 cl_int error = clGetDeviceInfo(
49 device, CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR,
50 sizeof(required_properties), &required_properties, NULL);
51 test_error(error,
52 "Unable to query "
53 "CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR");
54
55 cl_command_queue_properties queue_properties;
56 error = clGetCommandQueueInfo(queue, CL_QUEUE_PROPERTIES,
57 sizeof(queue_properties), &queue_properties,
58 NULL);
59 test_error(error, "Unable to query CL_QUEUE_PROPERTIES");
60
61
62 // Query if device supports simultaneous use
63 cl_device_command_buffer_capabilities_khr capabilities;
64 error = clGetDeviceInfo(device, CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR,
65 sizeof(capabilities), &capabilities, NULL);
66 test_error(error,
67 "Unable to query CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR");
68 simultaneous_use_support = simultaneous_use_requested
69 && (capabilities & CL_COMMAND_BUFFER_CAPABILITY_SIMULTANEOUS_USE_KHR)
70 != 0;
71 out_of_order_support =
72 capabilities & CL_COMMAND_BUFFER_CAPABILITY_OUT_OF_ORDER_KHR;
73
74 // Skip if queue properties don't contain those required
75 return required_properties != (required_properties & queue_properties);
76 }
77
78 //--------------------------------------------------------------------------
SetUpKernel()79 cl_int BasicCommandBufferTest::SetUpKernel()
80 {
81 cl_int error = CL_SUCCESS;
82
83 // Kernel performs a parallel copy from an input buffer to output buffer
84 // is created.
85 const char *kernel_str =
86 R"(
87 __kernel void copy(__global int* in, __global int* out, __global int* offset) {
88 size_t id = get_global_id(0);
89 int ind = offset[0] + id;
90 out[ind] = in[ind];
91 })";
92
93 error = create_single_kernel_helper_create_program(context, &program, 1,
94 &kernel_str);
95 test_error(error, "Failed to create program with source");
96
97 error = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
98 test_error(error, "Failed to build program");
99
100 kernel = clCreateKernel(program, "copy", &error);
101 test_error(error, "Failed to create copy kernel");
102
103 return CL_SUCCESS;
104 }
105
106 //--------------------------------------------------------------------------
SetUpKernelArgs()107 cl_int BasicCommandBufferTest::SetUpKernelArgs()
108 {
109 cl_int error = CL_SUCCESS;
110 in_mem =
111 clCreateBuffer(context, CL_MEM_READ_ONLY,
112 sizeof(cl_int) * num_elements * buffer_size_multiplier,
113 nullptr, &error);
114 test_error(error, "clCreateBuffer failed");
115
116 out_mem =
117 clCreateBuffer(context, CL_MEM_WRITE_ONLY,
118 sizeof(cl_int) * num_elements * buffer_size_multiplier,
119 nullptr, &error);
120 test_error(error, "clCreateBuffer failed");
121
122 cl_int offset = 0;
123 off_mem = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
124 sizeof(cl_int), &offset, &error);
125 test_error(error, "clCreateBuffer failed");
126
127 error = clSetKernelArg(kernel, 0, sizeof(in_mem), &in_mem);
128 test_error(error, "clSetKernelArg failed");
129
130 error = clSetKernelArg(kernel, 1, sizeof(out_mem), &out_mem);
131 test_error(error, "clSetKernelArg failed");
132
133 error = clSetKernelArg(kernel, 2, sizeof(off_mem), &off_mem);
134 test_error(error, "clSetKernelArg failed");
135
136 return CL_SUCCESS;
137 }
138
139 //--------------------------------------------------------------------------
SetUp(int elements)140 cl_int BasicCommandBufferTest::SetUp(int elements)
141 {
142 cl_int error = init_extension_functions();
143 if (error != CL_SUCCESS)
144 {
145 return error;
146 }
147
148 if (elements <= 0)
149 {
150 return CL_INVALID_VALUE;
151 }
152 num_elements = static_cast<size_t>(elements);
153
154 error = SetUpKernel();
155 test_error(error, "SetUpKernel failed");
156
157 error = SetUpKernelArgs();
158 test_error(error, "SetUpKernelArgs failed");
159
160 if (simultaneous_use_support)
161 {
162 cl_command_buffer_properties_khr properties[3] = {
163 CL_COMMAND_BUFFER_FLAGS_KHR, CL_COMMAND_BUFFER_SIMULTANEOUS_USE_KHR,
164 0
165 };
166 command_buffer =
167 clCreateCommandBufferKHR(1, &queue, properties, &error);
168 }
169 else
170 {
171 command_buffer = clCreateCommandBufferKHR(1, &queue, nullptr, &error);
172 }
173 test_error(error, "clCreateCommandBufferKHR failed");
174
175 return CL_SUCCESS;
176 }
177
178 namespace {
179
180 // Test enqueuing a command-buffer containing a single NDRange command once
181 struct BasicEnqueueTest : public BasicCommandBufferTest
182 {
183 using BasicCommandBufferTest::BasicCommandBufferTest;
184
Run__anon0db5bc270111::BasicEnqueueTest185 cl_int Run() override
186 {
187 cl_int error = clCommandNDRangeKernelKHR(
188 command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements,
189 nullptr, 0, nullptr, nullptr, nullptr);
190 test_error(error, "clCommandNDRangeKernelKHR failed");
191
192 error = clFinalizeCommandBufferKHR(command_buffer);
193 test_error(error, "clFinalizeCommandBufferKHR failed");
194
195 const cl_int pattern = 42;
196 error = clEnqueueFillBuffer(queue, in_mem, &pattern, sizeof(cl_int), 0,
197 data_size(), 0, nullptr, nullptr);
198 test_error(error, "clEnqueueFillBuffer failed");
199
200 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
201 nullptr, nullptr);
202 test_error(error, "clEnqueueCommandBufferKHR failed");
203
204 std::vector<cl_int> output_data_1(num_elements);
205 error = clEnqueueReadBuffer(queue, out_mem, CL_TRUE, 0, data_size(),
206 output_data_1.data(), 0, nullptr, nullptr);
207 test_error(error, "clEnqueueReadBuffer failed");
208
209 for (size_t i = 0; i < num_elements; i++)
210 {
211 CHECK_VERIFICATION_ERROR(pattern, output_data_1[i], i);
212 }
213
214 const cl_int new_pattern = 12;
215 error = clEnqueueFillBuffer(queue, in_mem, &new_pattern, sizeof(cl_int),
216 0, data_size(), 0, nullptr, nullptr);
217 test_error(error, "clEnqueueFillBuffer failed");
218
219 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
220 nullptr, nullptr);
221 test_error(error, "clEnqueueCommandBufferKHR failed");
222
223 std::vector<cl_int> output_data_2(num_elements);
224 error = clEnqueueReadBuffer(queue, out_mem, CL_TRUE, 0, data_size(),
225 output_data_2.data(), 0, nullptr, nullptr);
226 test_error(error, "clEnqueueReadBuffer failed");
227
228 for (size_t i = 0; i < num_elements; i++)
229 {
230 CHECK_VERIFICATION_ERROR(new_pattern, output_data_2[i], i);
231 }
232
233 return CL_SUCCESS;
234 }
235 };
236
237 // Test enqueuing a command-buffer containing multiple command, including
238 // operations other than NDRange kernel execution.
239 struct MixedCommandsTest : public BasicCommandBufferTest
240 {
241 using BasicCommandBufferTest::BasicCommandBufferTest;
242
Run__anon0db5bc270111::MixedCommandsTest243 cl_int Run() override
244 {
245 cl_int error;
246 const size_t iterations = 4;
247 clMemWrapper result_mem =
248 clCreateBuffer(context, CL_MEM_READ_WRITE,
249 sizeof(cl_int) * iterations, nullptr, &error);
250 test_error(error, "clCreateBuffer failed");
251
252 const cl_int pattern_base = 42;
253 for (size_t i = 0; i < iterations; i++)
254 {
255 const cl_int pattern = pattern_base + i;
256 cl_int error = clCommandFillBufferKHR(
257 command_buffer, nullptr, in_mem, &pattern, sizeof(cl_int), 0,
258 data_size(), 0, nullptr, nullptr, nullptr);
259 test_error(error, "clCommandFillBufferKHR failed");
260
261 error = clCommandNDRangeKernelKHR(
262 command_buffer, nullptr, nullptr, kernel, 1, nullptr,
263 &num_elements, nullptr, 0, nullptr, nullptr, nullptr);
264 test_error(error, "clCommandNDRangeKernelKHR failed");
265
266 const size_t result_offset = i * sizeof(cl_int);
267 error = clCommandCopyBufferKHR(
268 command_buffer, nullptr, out_mem, result_mem, 0, result_offset,
269 sizeof(cl_int), 0, nullptr, nullptr, nullptr);
270 test_error(error, "clCommandCopyBufferKHR failed");
271 }
272
273 error = clFinalizeCommandBufferKHR(command_buffer);
274 test_error(error, "clFinalizeCommandBufferKHR failed");
275
276 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
277 nullptr, nullptr);
278 test_error(error, "clEnqueueCommandBufferKHR failed");
279
280 std::vector<cl_int> result_data(num_elements);
281 error = clEnqueueReadBuffer(queue, result_mem, CL_TRUE, 0,
282 iterations * sizeof(cl_int),
283 result_data.data(), 0, nullptr, nullptr);
284 test_error(error, "clEnqueueReadBuffer failed");
285
286 for (size_t i = 0; i < iterations; i++)
287 {
288 const cl_int ref = pattern_base + i;
289 CHECK_VERIFICATION_ERROR(ref, result_data[i], i);
290 }
291
292 return CL_SUCCESS;
293 }
294 };
295
296 // Test flushing the command-queue between command-buffer enqueues
297 struct ExplicitFlushTest : public BasicCommandBufferTest
298 {
299 using BasicCommandBufferTest::BasicCommandBufferTest;
300
Run__anon0db5bc270111::ExplicitFlushTest301 cl_int Run() override
302 {
303 cl_int error = clCommandNDRangeKernelKHR(
304 command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements,
305 nullptr, 0, nullptr, nullptr, nullptr);
306 test_error(error, "clCommandNDRangeKernelKHR failed");
307
308 error = clFinalizeCommandBufferKHR(command_buffer);
309 test_error(error, "clFinalizeCommandBufferKHR failed");
310
311 const cl_int pattern_A = 42;
312 error = clEnqueueFillBuffer(queue, in_mem, &pattern_A, sizeof(cl_int),
313 0, data_size(), 0, nullptr, nullptr);
314 test_error(error, "clEnqueueFillBuffer failed");
315
316 error = clFlush(queue);
317 test_error(error, "clFlush failed");
318
319 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
320 nullptr, nullptr);
321 test_error(error, "clEnqueueCommandBufferKHR failed");
322
323 std::vector<cl_int> output_data_A(num_elements);
324 error = clEnqueueReadBuffer(queue, out_mem, CL_FALSE, 0, data_size(),
325 output_data_A.data(), 0, nullptr, nullptr);
326 test_error(error, "clEnqueueReadBuffer failed");
327
328 const cl_int pattern_B = 0xA;
329 error = clEnqueueFillBuffer(queue, in_mem, &pattern_B, sizeof(cl_int),
330 0, data_size(), 0, nullptr, nullptr);
331 test_error(error, "clEnqueueFillBuffer failed");
332
333 error = clFlush(queue);
334 test_error(error, "clFlush failed");
335
336 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
337 nullptr, nullptr);
338 test_error(error, "clEnqueueCommandBufferKHR failed");
339
340 error = clFlush(queue);
341 test_error(error, "clFlush failed");
342
343 std::vector<cl_int> output_data_B(num_elements);
344 error = clEnqueueReadBuffer(queue, out_mem, CL_FALSE, 0, data_size(),
345 output_data_B.data(), 0, nullptr, nullptr);
346 test_error(error, "clEnqueueReadBuffer failed");
347
348 error = clFinish(queue);
349 test_error(error, "clFinish failed");
350
351 for (size_t i = 0; i < num_elements; i++)
352 {
353 CHECK_VERIFICATION_ERROR(pattern_A, output_data_A[i], i);
354
355 CHECK_VERIFICATION_ERROR(pattern_B, output_data_B[i], i);
356 }
357 return CL_SUCCESS;
358 }
359
Skip__anon0db5bc270111::ExplicitFlushTest360 bool Skip() override
361 {
362 return BasicCommandBufferTest::Skip() || !simultaneous_use_support;
363 }
364 };
365
366 // Test enqueueing a command-buffer twice separated by another enqueue operation
367 struct InterleavedEnqueueTest : public BasicCommandBufferTest
368 {
369 using BasicCommandBufferTest::BasicCommandBufferTest;
370
Run__anon0db5bc270111::InterleavedEnqueueTest371 cl_int Run() override
372 {
373 cl_int error = clCommandNDRangeKernelKHR(
374 command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements,
375 nullptr, 0, nullptr, nullptr, nullptr);
376 test_error(error, "clCommandNDRangeKernelKHR failed");
377
378 error = clFinalizeCommandBufferKHR(command_buffer);
379 test_error(error, "clFinalizeCommandBufferKHR failed");
380
381 cl_int pattern = 42;
382 error = clEnqueueFillBuffer(queue, in_mem, &pattern, sizeof(cl_int), 0,
383 data_size(), 0, nullptr, nullptr);
384 test_error(error, "clEnqueueFillBuffer failed");
385
386 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
387 nullptr, nullptr);
388 test_error(error, "clEnqueueCommandBufferKHR failed");
389
390 pattern = 0xABCD;
391 error = clEnqueueFillBuffer(queue, in_mem, &pattern, sizeof(cl_int), 0,
392 data_size(), 0, nullptr, nullptr);
393 test_error(error, "clEnqueueFillBuffer failed");
394
395 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
396 nullptr, nullptr);
397 test_error(error, "clEnqueueCommandBufferKHR failed");
398
399 error = clEnqueueCopyBuffer(queue, in_mem, out_mem, 0, 0, data_size(),
400 0, nullptr, nullptr);
401 test_error(error, "clEnqueueCopyBuffer failed");
402
403 std::vector<cl_int> output_data(num_elements);
404 error = clEnqueueReadBuffer(queue, out_mem, CL_TRUE, 0, data_size(),
405 output_data.data(), 0, nullptr, nullptr);
406 test_error(error, "clEnqueueReadBuffer failed");
407
408 for (size_t i = 0; i < num_elements; i++)
409 {
410 CHECK_VERIFICATION_ERROR(pattern, output_data[i], i);
411 }
412
413 return CL_SUCCESS;
414 }
415
Skip__anon0db5bc270111::InterleavedEnqueueTest416 bool Skip() override
417 {
418 return BasicCommandBufferTest::Skip() || !simultaneous_use_support;
419 }
420 };
421
422 } // anonymous namespace
423
test_single_ndrange(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)424 int test_single_ndrange(cl_device_id device, cl_context context,
425 cl_command_queue queue, int num_elements)
426 {
427 return MakeAndRunTest<BasicEnqueueTest>(device, context, queue,
428 num_elements);
429 }
430
test_interleaved_enqueue(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)431 int test_interleaved_enqueue(cl_device_id device, cl_context context,
432 cl_command_queue queue, int num_elements)
433 {
434 return MakeAndRunTest<InterleavedEnqueueTest>(device, context, queue,
435 num_elements);
436 }
437
test_mixed_commands(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)438 int test_mixed_commands(cl_device_id device, cl_context context,
439 cl_command_queue queue, int num_elements)
440 {
441 return MakeAndRunTest<MixedCommandsTest>(device, context, queue,
442 num_elements);
443 }
444
test_explicit_flush(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)445 int test_explicit_flush(cl_device_id device, cl_context context,
446 cl_command_queue queue, int num_elements)
447 {
448 return MakeAndRunTest<ExplicitFlushTest>(device, context, queue,
449 num_elements);
450 }
451