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 "command_buffer_test_base.h"
17 #include "procs.h"
18 #include "harness/typeWrappers.h"
19
20 #include <algorithm>
21 #include <cstring>
22 #include <vector>
23
24 #define CHECK_VERIFICATION_ERROR(reference, result, index) \
25 { \
26 if (reference != result) \
27 { \
28 log_error("Expected %d was %d at index %u\n", reference, result, \
29 index); \
30 return TEST_FAIL; \
31 } \
32 }
33
34 namespace {
35
36 // Helper test fixture for constructing OpenCL objects used in testing
37 // a variety of simple command-buffer enqueue scenarios.
38 struct BasicCommandBufferTest : CommandBufferTestBase
39 {
40
BasicCommandBufferTest__anon1a84e5d80111::BasicCommandBufferTest41 BasicCommandBufferTest(cl_device_id device, cl_context context,
42 cl_command_queue queue)
43 : CommandBufferTestBase(device), context(context), queue(queue),
44 command_buffer(this), simultaneous_use(false),
45 out_of_order_support(false), num_elements(0)
46 {}
47
Skip__anon1a84e5d80111::BasicCommandBufferTest48 virtual bool Skip()
49 {
50 cl_command_queue_properties required_properties;
51 cl_int error = clGetDeviceInfo(
52 device, CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR,
53 sizeof(required_properties), &required_properties, NULL);
54 test_error(error,
55 "Unable to query "
56 "CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR");
57
58 cl_command_queue_properties queue_properties;
59
60 error = clGetCommandQueueInfo(queue, CL_QUEUE_PROPERTIES,
61 sizeof(queue_properties),
62 &queue_properties, NULL);
63 test_error(error, "Unable to query CL_QUEUE_PROPERTIES");
64
65 // Skip if queue properties don't contain those required
66 return required_properties != (required_properties & queue_properties);
67 }
68
SetUp__anon1a84e5d80111::BasicCommandBufferTest69 virtual cl_int SetUp(int elements)
70 {
71 cl_int error = init_extension_functions();
72 if (error != CL_SUCCESS)
73 {
74 return error;
75 }
76
77 // Query if device supports simultaneous use
78 cl_device_command_buffer_capabilities_khr capabilities;
79 error =
80 clGetDeviceInfo(device, CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR,
81 sizeof(capabilities), &capabilities, NULL);
82 test_error(error,
83 "Unable to query CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR");
84 simultaneous_use =
85 capabilities & CL_COMMAND_BUFFER_CAPABILITY_SIMULTANEOUS_USE_KHR;
86 out_of_order_support =
87 capabilities & CL_COMMAND_BUFFER_CAPABILITY_OUT_OF_ORDER_KHR;
88
89 if (elements <= 0)
90 {
91 return CL_INVALID_VALUE;
92 }
93 num_elements = static_cast<size_t>(elements);
94
95 // Kernel performs a parallel copy from an input buffer to output buffer
96 // is created.
97 const char *kernel_str =
98 R"(
99 __kernel void copy(__global int* in, __global int* out) {
100 size_t id = get_global_id(0);
101 out[id] = in[id];
102 })";
103
104 error = create_single_kernel_helper_create_program(context, &program, 1,
105 &kernel_str);
106 test_error(error, "Failed to create program with source");
107
108 error = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
109 test_error(error, "Failed to build program");
110
111 in_mem = clCreateBuffer(context, CL_MEM_READ_ONLY,
112 sizeof(cl_int) * num_elements, nullptr, &error);
113 test_error(error, "clCreateBuffer failed");
114
115 out_mem =
116 clCreateBuffer(context, CL_MEM_WRITE_ONLY,
117 sizeof(cl_int) * num_elements, nullptr, &error);
118 test_error(error, "clCreateBuffer failed");
119
120 kernel = clCreateKernel(program, "copy", &error);
121 test_error(error, "Failed to create copy kernel");
122
123 error = clSetKernelArg(kernel, 0, sizeof(in_mem), &in_mem);
124 test_error(error, "clSetKernelArg failed");
125
126 error = clSetKernelArg(kernel, 1, sizeof(out_mem), &out_mem);
127 test_error(error, "clSetKernelArg failed");
128
129 if (simultaneous_use)
130 {
131 cl_command_buffer_properties_khr properties[3] = {
132 CL_COMMAND_BUFFER_FLAGS_KHR,
133 CL_COMMAND_BUFFER_SIMULTANEOUS_USE_KHR, 0
134 };
135 command_buffer =
136 clCreateCommandBufferKHR(1, &queue, properties, &error);
137 }
138 else
139 {
140 command_buffer =
141 clCreateCommandBufferKHR(1, &queue, nullptr, &error);
142 }
143 test_error(error, "clCreateCommandBufferKHR failed");
144
145 return CL_SUCCESS;
146 }
147
148 // Test body returning an OpenCL error code
149 virtual cl_int Run() = 0;
150
151
152 protected:
data_size__anon1a84e5d80111::BasicCommandBufferTest153 size_t data_size() const { return num_elements * sizeof(cl_int); }
154
155 cl_context context;
156 cl_command_queue queue;
157 clCommandBufferWrapper command_buffer;
158 clProgramWrapper program;
159 clKernelWrapper kernel;
160 clMemWrapper in_mem, out_mem;
161 size_t num_elements;
162
163 // Device support query results
164 bool simultaneous_use;
165 bool out_of_order_support;
166 };
167
168 // Test enqueuing a command-buffer containing a single NDRange command once
169 struct BasicEnqueueTest : public BasicCommandBufferTest
170 {
171 using BasicCommandBufferTest::BasicCommandBufferTest;
172
Run__anon1a84e5d80111::BasicEnqueueTest173 cl_int Run() override
174 {
175 cl_int error = clCommandNDRangeKernelKHR(
176 command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements,
177 nullptr, 0, nullptr, nullptr, nullptr);
178 test_error(error, "clCommandNDRangeKernelKHR failed");
179
180 error = clFinalizeCommandBufferKHR(command_buffer);
181 test_error(error, "clFinalizeCommandBufferKHR failed");
182
183 const cl_int pattern = 42;
184 error = clEnqueueFillBuffer(queue, in_mem, &pattern, sizeof(cl_int), 0,
185 data_size(), 0, nullptr, nullptr);
186 test_error(error, "clEnqueueFillBuffer failed");
187
188 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
189 nullptr, nullptr);
190 test_error(error, "clEnqueueCommandBufferKHR failed");
191
192 std::vector<cl_int> output_data(num_elements);
193 error = clEnqueueReadBuffer(queue, out_mem, CL_TRUE, 0, data_size(),
194 output_data.data(), 0, nullptr, nullptr);
195 test_error(error, "clEnqueueReadBuffer failed");
196
197 for (size_t i = 0; i < num_elements; i++)
198 {
199 CHECK_VERIFICATION_ERROR(pattern, output_data[i], i);
200 }
201
202 return CL_SUCCESS;
203 }
204 };
205
206 // Test enqueuing a command-buffer containing multiple command, including
207 // operations other than NDRange kernel execution.
208 struct MixedCommandsTest : public BasicCommandBufferTest
209 {
210 using BasicCommandBufferTest::BasicCommandBufferTest;
211
Run__anon1a84e5d80111::MixedCommandsTest212 cl_int Run() override
213 {
214 cl_int error;
215 const size_t iterations = 4;
216 clMemWrapper result_mem =
217 clCreateBuffer(context, CL_MEM_READ_WRITE,
218 sizeof(cl_int) * iterations, nullptr, &error);
219 test_error(error, "clCreateBuffer failed");
220
221 const cl_int pattern_base = 42;
222 for (size_t i = 0; i < iterations; i++)
223 {
224 const cl_int pattern = pattern_base + i;
225 cl_int error = clCommandFillBufferKHR(
226 command_buffer, nullptr, in_mem, &pattern, sizeof(cl_int), 0,
227 data_size(), 0, nullptr, nullptr, nullptr);
228 test_error(error, "clCommandFillBufferKHR failed");
229
230 error = clCommandNDRangeKernelKHR(
231 command_buffer, nullptr, nullptr, kernel, 1, nullptr,
232 &num_elements, nullptr, 0, nullptr, nullptr, nullptr);
233 test_error(error, "clCommandNDRangeKernelKHR failed");
234
235 const size_t result_offset = i * sizeof(cl_int);
236 error = clCommandCopyBufferKHR(
237 command_buffer, nullptr, out_mem, result_mem, 0, result_offset,
238 sizeof(cl_int), 0, nullptr, nullptr, nullptr);
239 test_error(error, "clCommandCopyBufferKHR failed");
240 }
241
242 error = clFinalizeCommandBufferKHR(command_buffer);
243 test_error(error, "clFinalizeCommandBufferKHR failed");
244
245 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
246 nullptr, nullptr);
247 test_error(error, "clEnqueueCommandBufferKHR failed");
248
249 std::vector<cl_int> result_data(num_elements);
250 error = clEnqueueReadBuffer(queue, result_mem, CL_TRUE, 0,
251 iterations * sizeof(cl_int),
252 result_data.data(), 0, nullptr, nullptr);
253 test_error(error, "clEnqueueReadBuffer failed");
254
255 for (size_t i = 0; i < iterations; i++)
256 {
257 const cl_int ref = pattern_base + i;
258 CHECK_VERIFICATION_ERROR(ref, result_data[i], i);
259 }
260
261 return CL_SUCCESS;
262 }
263 };
264
265 // Test enqueueing a command-buffer blocked on a user-event
266 struct UserEventTest : public BasicCommandBufferTest
267 {
268 using BasicCommandBufferTest::BasicCommandBufferTest;
269
Run__anon1a84e5d80111::UserEventTest270 cl_int Run() override
271 {
272 cl_int error = clCommandNDRangeKernelKHR(
273 command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements,
274 nullptr, 0, nullptr, nullptr, nullptr);
275 test_error(error, "clCommandNDRangeKernelKHR failed");
276
277 error = clFinalizeCommandBufferKHR(command_buffer);
278 test_error(error, "clFinalizeCommandBufferKHR failed");
279
280 clEventWrapper user_event = clCreateUserEvent(context, &error);
281 test_error(error, "clCreateUserEvent failed");
282
283 const cl_int pattern = 42;
284 error = clEnqueueFillBuffer(queue, in_mem, &pattern, sizeof(cl_int), 0,
285 data_size(), 0, nullptr, nullptr);
286 test_error(error, "clEnqueueFillBuffer failed");
287
288 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 1,
289 &user_event, nullptr);
290 test_error(error, "clEnqueueCommandBufferKHR failed");
291
292 std::vector<cl_int> output_data(num_elements);
293 error = clEnqueueReadBuffer(queue, out_mem, CL_FALSE, 0, data_size(),
294 output_data.data(), 0, nullptr, nullptr);
295 test_error(error, "clEnqueueReadBuffer failed");
296
297 error = clSetUserEventStatus(user_event, CL_COMPLETE);
298 test_error(error, "clSetUserEventStatus failed");
299
300 error = clFinish(queue);
301 test_error(error, "clFinish failed");
302
303 for (size_t i = 0; i < num_elements; i++)
304 {
305 CHECK_VERIFICATION_ERROR(pattern, output_data[i], i);
306 }
307
308 return CL_SUCCESS;
309 }
310 };
311
312 // Test flushing the command-queue between command-buffer enqueues
313 struct ExplicitFlushTest : public BasicCommandBufferTest
314 {
315 using BasicCommandBufferTest::BasicCommandBufferTest;
316
Run__anon1a84e5d80111::ExplicitFlushTest317 cl_int Run() override
318 {
319 cl_int error = clCommandNDRangeKernelKHR(
320 command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements,
321 nullptr, 0, nullptr, nullptr, nullptr);
322 test_error(error, "clCommandNDRangeKernelKHR failed");
323
324 error = clFinalizeCommandBufferKHR(command_buffer);
325 test_error(error, "clFinalizeCommandBufferKHR failed");
326
327 const cl_int pattern_A = 42;
328 error = clEnqueueFillBuffer(queue, in_mem, &pattern_A, sizeof(cl_int),
329 0, data_size(), 0, nullptr, nullptr);
330 test_error(error, "clEnqueueFillBuffer failed");
331
332 error = clFlush(queue);
333 test_error(error, "clFlush failed");
334
335 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
336 nullptr, nullptr);
337 test_error(error, "clEnqueueCommandBufferKHR failed");
338
339 std::vector<cl_int> output_data_A(num_elements);
340 error = clEnqueueReadBuffer(queue, out_mem, CL_FALSE, 0, data_size(),
341 output_data_A.data(), 0, nullptr, nullptr);
342 test_error(error, "clEnqueueReadBuffer failed");
343
344 const cl_int pattern_B = 0xA;
345 error = clEnqueueFillBuffer(queue, in_mem, &pattern_B, sizeof(cl_int),
346 0, data_size(), 0, nullptr, nullptr);
347 test_error(error, "clEnqueueFillBuffer failed");
348
349 error = clFlush(queue);
350 test_error(error, "clFlush failed");
351
352 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
353 nullptr, nullptr);
354 test_error(error, "clEnqueueCommandBufferKHR failed");
355
356 error = clFlush(queue);
357 test_error(error, "clFlush failed");
358
359 std::vector<cl_int> output_data_B(num_elements);
360 error = clEnqueueReadBuffer(queue, out_mem, CL_FALSE, 0, data_size(),
361 output_data_B.data(), 0, nullptr, nullptr);
362 test_error(error, "clEnqueueReadBuffer failed");
363
364 error = clFinish(queue);
365 test_error(error, "clFinish failed");
366
367 for (size_t i = 0; i < num_elements; i++)
368 {
369 CHECK_VERIFICATION_ERROR(pattern_A, output_data_A[i], i);
370
371 CHECK_VERIFICATION_ERROR(pattern_B, output_data_B[i], i);
372 }
373 return CL_SUCCESS;
374 }
375
Skip__anon1a84e5d80111::ExplicitFlushTest376 bool Skip() override
377 {
378 return !simultaneous_use || BasicCommandBufferTest::Skip();
379 }
380 };
381
382 // Test enqueueing a command-buffer twice separated by another enqueue operation
383 struct InterleavedEnqueueTest : public BasicCommandBufferTest
384 {
385 using BasicCommandBufferTest::BasicCommandBufferTest;
386
Run__anon1a84e5d80111::InterleavedEnqueueTest387 cl_int Run() override
388 {
389 cl_int error = clCommandNDRangeKernelKHR(
390 command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements,
391 nullptr, 0, nullptr, nullptr, nullptr);
392 test_error(error, "clCommandNDRangeKernelKHR failed");
393
394 error = clFinalizeCommandBufferKHR(command_buffer);
395 test_error(error, "clFinalizeCommandBufferKHR failed");
396
397 cl_int pattern = 42;
398 error = clEnqueueFillBuffer(queue, in_mem, &pattern, sizeof(cl_int), 0,
399 data_size(), 0, nullptr, nullptr);
400 test_error(error, "clEnqueueFillBuffer failed");
401
402 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
403 nullptr, nullptr);
404 test_error(error, "clEnqueueCommandBufferKHR failed");
405
406 pattern = 0xABCD;
407 error = clEnqueueFillBuffer(queue, in_mem, &pattern, sizeof(cl_int), 0,
408 data_size(), 0, nullptr, nullptr);
409 test_error(error, "clEnqueueFillBuffer failed");
410
411 error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
412 nullptr, nullptr);
413 test_error(error, "clEnqueueCommandBufferKHR failed");
414
415 error = clEnqueueCopyBuffer(queue, in_mem, out_mem, 0, 0, data_size(),
416 0, nullptr, nullptr);
417 test_error(error, "clEnqueueCopyBuffer failed");
418
419 std::vector<cl_int> output_data(num_elements);
420 error = clEnqueueReadBuffer(queue, out_mem, CL_TRUE, 0, data_size(),
421 output_data.data(), 0, nullptr, nullptr);
422 test_error(error, "clEnqueueReadBuffer failed");
423
424 for (size_t i = 0; i < num_elements; i++)
425 {
426 CHECK_VERIFICATION_ERROR(pattern, output_data[i], i);
427 }
428
429 return CL_SUCCESS;
430 }
431
Skip__anon1a84e5d80111::InterleavedEnqueueTest432 bool Skip() override
433 {
434 return !simultaneous_use || BasicCommandBufferTest::Skip();
435 }
436 };
437
438 // Test sync-points with an out-of-order command-buffer
439 struct OutOfOrderTest : public BasicCommandBufferTest
440 {
441 using BasicCommandBufferTest::BasicCommandBufferTest;
OutOfOrderTest__anon1a84e5d80111::OutOfOrderTest442 OutOfOrderTest(cl_device_id device, cl_context context,
443 cl_command_queue queue)
444 : BasicCommandBufferTest(device, context, queue),
445 out_of_order_command_buffer(this), out_of_order_queue(nullptr),
446 event(nullptr)
447 {}
448
Run__anon1a84e5d80111::OutOfOrderTest449 cl_int Run() override
450 {
451 cl_sync_point_khr sync_points[2];
452
453 const cl_int pattern = 42;
454 cl_int error =
455 clCommandFillBufferKHR(out_of_order_command_buffer, nullptr, in_mem,
456 &pattern, sizeof(cl_int), 0, data_size(), 0,
457 nullptr, &sync_points[0], nullptr);
458 test_error(error, "clCommandFillBufferKHR failed");
459
460 const cl_int overwritten_pattern = 0xACDC;
461 error = clCommandFillBufferKHR(out_of_order_command_buffer, nullptr,
462 out_mem, &overwritten_pattern,
463 sizeof(cl_int), 0, data_size(), 0,
464 nullptr, &sync_points[1], nullptr);
465 test_error(error, "clCommandFillBufferKHR failed");
466
467 error = clCommandNDRangeKernelKHR(
468 out_of_order_command_buffer, nullptr, nullptr, kernel, 1, nullptr,
469 &num_elements, nullptr, 2, sync_points, nullptr, nullptr);
470 test_error(error, "clCommandNDRangeKernelKHR failed");
471
472 error = clFinalizeCommandBufferKHR(out_of_order_command_buffer);
473 test_error(error, "clFinalizeCommandBufferKHR failed");
474
475 error = clEnqueueCommandBufferKHR(
476 0, nullptr, out_of_order_command_buffer, 0, nullptr, &event);
477 test_error(error, "clEnqueueCommandBufferKHR failed");
478
479 std::vector<cl_int> output_data(num_elements);
480 error = clEnqueueReadBuffer(out_of_order_queue, out_mem, CL_TRUE, 0,
481 data_size(), output_data.data(), 1, &event,
482 nullptr);
483 test_error(error, "clEnqueueReadBuffer failed");
484
485 for (size_t i = 0; i < num_elements; i++)
486 {
487 CHECK_VERIFICATION_ERROR(pattern, output_data[i], i);
488 }
489
490 return CL_SUCCESS;
491 }
492
SetUp__anon1a84e5d80111::OutOfOrderTest493 cl_int SetUp(int elements) override
494 {
495 cl_int error = BasicCommandBufferTest::SetUp(elements);
496 test_error(error, "BasicCommandBufferTest::SetUp failed");
497
498 if (!out_of_order_support)
499 {
500 // Test will skip as device doesn't support out-of-order
501 // command-buffers
502 return CL_SUCCESS;
503 }
504
505 out_of_order_queue = clCreateCommandQueue(
506 context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &error);
507 test_error(error, "Unable to create command queue to test with");
508
509 out_of_order_command_buffer =
510 clCreateCommandBufferKHR(1, &out_of_order_queue, nullptr, &error);
511 test_error(error, "clCreateCommandBufferKHR failed");
512
513 return CL_SUCCESS;
514 }
515
Skip__anon1a84e5d80111::OutOfOrderTest516 bool Skip() override
517 {
518 return !out_of_order_support || BasicCommandBufferTest::Skip();
519 }
520
521 clCommandQueueWrapper out_of_order_queue;
522 clCommandBufferWrapper out_of_order_command_buffer;
523 clEventWrapper event;
524 };
525
526 #undef CHECK_VERIFICATION_ERROR
527
528 template <class T>
MakeAndRunTest(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)529 int MakeAndRunTest(cl_device_id device, cl_context context,
530 cl_command_queue queue, int num_elements)
531 {
532 CHECK_COMMAND_BUFFER_EXTENSION_AVAILABLE(device);
533
534 auto test_fixture = T(device, context, queue);
535 cl_int error = test_fixture.SetUp(num_elements);
536 test_error_ret(error, "Error in test initialization", TEST_FAIL);
537
538 if (test_fixture.Skip())
539 {
540 return TEST_SKIPPED_ITSELF;
541 }
542
543 error = test_fixture.Run();
544 test_error_ret(error, "Test Failed", TEST_FAIL);
545
546 return TEST_PASS;
547 }
548 } // anonymous namespace
549
test_single_ndrange(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)550 int test_single_ndrange(cl_device_id device, cl_context context,
551 cl_command_queue queue, int num_elements)
552 {
553 return MakeAndRunTest<BasicEnqueueTest>(device, context, queue,
554 num_elements);
555 }
556
test_interleaved_enqueue(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)557 int test_interleaved_enqueue(cl_device_id device, cl_context context,
558 cl_command_queue queue, int num_elements)
559 {
560 return MakeAndRunTest<InterleavedEnqueueTest>(device, context, queue,
561 num_elements);
562 }
563
test_mixed_commands(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)564 int test_mixed_commands(cl_device_id device, cl_context context,
565 cl_command_queue queue, int num_elements)
566 {
567 return MakeAndRunTest<MixedCommandsTest>(device, context, queue,
568 num_elements);
569 }
570
test_explicit_flush(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)571 int test_explicit_flush(cl_device_id device, cl_context context,
572 cl_command_queue queue, int num_elements)
573 {
574 return MakeAndRunTest<ExplicitFlushTest>(device, context, queue,
575 num_elements);
576 }
577
test_user_events(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)578 int test_user_events(cl_device_id device, cl_context context,
579 cl_command_queue queue, int num_elements)
580 {
581 return MakeAndRunTest<UserEventTest>(device, context, queue, num_elements);
582 }
583
test_out_of_order(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)584 int test_out_of_order(cl_device_id device, cl_context context,
585 cl_command_queue queue, int num_elements)
586 {
587 return MakeAndRunTest<OutOfOrderTest>(device, context, queue, num_elements);
588 }
589