Initial command-buffer extension tests (#1368)

* Initial command-buffer tests

Introduce some basic testing of the
[cl_khr_command_buffer](https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_Ext.html#cl_khr_command_buffer)
extension. This is intended as a starting point from which we can iteratively build up tests
for the extension collaboratively.

* Move tests into derived classes

* Move tests from methods into derived classes implementing
  a `Run()` interface.
* Fix memory leak when command_buffer isn't freed when a test
  is skipped.
* Print correct error code for
  `CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR`
* Pass `nullptr` for queue parameter to command recording entry-points

* Define command-buffer type wrapper

Other OpenCL object have a wrapper to reference count their use
and free the wrapped object. The command-buffer object can't use
the generic type wrappers which are templated on the appropriate
release/retain function, as the release/retain functions are
queried at runtime.

Instead, define our own command-buffer wrapper class where a base
object is passed on construction which contains function pointers
to the release/retain functions that can be used in the wrapper.

* Use create_single_kernel_helper_create_program

Use `create_single_kernel_helper_create_program` rather than
hardcoding `clCreateProgramWithSource` to allow for other types
of program input.

Also fix bug using wrong enum for passing properties on command-buffer
creation, should be `CL_COMMAND_BUFFER_FLAGS_KHR`

* Add out-of-order command-buffer test

Introduce a basic test for checking sync-point use
with out-of-order command-buffers.

This also includes better checking of required queue properties.

test_conformance/extensions/CMakeLists.txt

@@ -1,2 +1,3 @@
 add_subdirectory( cl_ext_cxx_for_opencl )
+add_subdirectory( cl_khr_command_buffer )
 add_subdirectory( cl_khr_dx9_media_sharing )

test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt

@@ -0,0 +1,8 @@
+set(MODULE_NAME CL_KHR_COMMAND_BUFFER)
+
+set(${MODULE_NAME}_SOURCES
+    main.cpp
+    basic_command_buffer.cpp
+)
+
+include(../../CMakeCommon.txt)

test_conformance/extensions/cl_khr_command_buffer/basic_command_buffer.cpp

@@ -0,0 +1,588 @@
+//
+// Copyright (c) 2022 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+#include "command_buffer_test_base.h"
+#include "procs.h"
+#include "harness/typeWrappers.h"
+
+#include <algorithm>
+#include <cstring>
+#include <vector>
+
+#define CHECK_VERIFICATION_ERROR(reference, result, index)                     \
+    {                                                                          \
+        if (reference != result)                                               \
+        {                                                                      \
+            log_error("Expected %d was %d at index %u\n", reference, result,   \
+                      index);                                                  \
+            return TEST_FAIL;                                                  \
+        }                                                                      \
+    }
+
+namespace {
+
+// Helper test fixture for constructing OpenCL objects used in testing
+// a variety of simple command-buffer enqueue scenarios.
+struct BasicCommandBufferTest : CommandBufferTestBase
+{
+
+    BasicCommandBufferTest(cl_device_id device, cl_context context,
+                           cl_command_queue queue)
+        : CommandBufferTestBase(device), context(context), queue(queue),
+          command_buffer(this), simultaneous_use(false),
+          out_of_order_support(false), num_elements(0)
+    {}
+
+    virtual bool Skip()
+    {
+        cl_command_queue_properties required_properties;
+        cl_int error = clGetDeviceInfo(
+            device, CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR,
+            sizeof(required_properties), &required_properties, NULL);
+        test_error(error,
+                   "Unable to query "
+                   "CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR");
+
+        cl_command_queue_properties queue_properties;
+
+        error = clGetCommandQueueInfo(queue, CL_QUEUE_PROPERTIES,
+                                      sizeof(queue_properties),
+                                      &queue_properties, NULL);
+        test_error(error, "Unable to query CL_QUEUE_PROPERTIES");
+
+        // Skip if queue properties don't contain those required
+        return required_properties != (required_properties & queue_properties);
+    }
+
+    virtual cl_int SetUp(int elements)
+    {
+        cl_int error = init_extension_functions();
+        if (error != CL_SUCCESS)
+        {
+            return error;
+        }
+
+        // Query if device supports simultaneous use
+        cl_device_command_buffer_capabilities_khr capabilities;
+        error =
+            clGetDeviceInfo(device, CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR,
+                            sizeof(capabilities), &capabilities, NULL);
+        test_error(error,
+                   "Unable to query CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR");
+        simultaneous_use =
+            capabilities & CL_COMMAND_BUFFER_CAPABILITY_SIMULTANEOUS_USE_KHR;
+        out_of_order_support =
+            capabilities & CL_COMMAND_BUFFER_CAPABILITY_OUT_OF_ORDER_KHR;
+
+        if (elements <= 0)
+        {
+            return CL_INVALID_VALUE;
+        }
+        num_elements = static_cast<size_t>(elements);
+
+        // Kernel performs a parallel copy from an input buffer to output buffer
+        // is created.
+        const char *kernel_str =
+            R"(
+        __kernel void copy(__global int* in, __global int* out) {
+            size_t id = get_global_id(0);
+            out[id] = in[id];
+        })";
+
+        error = create_single_kernel_helper_create_program(context, &program, 1,
+                                                           &kernel_str);
+        test_error(error, "Failed to create program with source");
+
+        error = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
+        test_error(error, "Failed to build program");
+
+        in_mem = clCreateBuffer(context, CL_MEM_READ_ONLY,
+                                sizeof(cl_int) * num_elements, nullptr, &error);
+        test_error(error, "clCreateBuffer failed");
+
+        out_mem =
+            clCreateBuffer(context, CL_MEM_WRITE_ONLY,
+                           sizeof(cl_int) * num_elements, nullptr, &error);
+        test_error(error, "clCreateBuffer failed");
+
+        kernel = clCreateKernel(program, "copy", &error);
+        test_error(error, "Failed to create copy kernel");
+
+        error = clSetKernelArg(kernel, 0, sizeof(in_mem), &in_mem);
+        test_error(error, "clSetKernelArg failed");
+
+        error = clSetKernelArg(kernel, 1, sizeof(out_mem), &out_mem);
+        test_error(error, "clSetKernelArg failed");
+
+        if (simultaneous_use)
+        {
+            cl_command_buffer_properties_khr properties[3] = {
+                CL_COMMAND_BUFFER_FLAGS_KHR,
+                CL_COMMAND_BUFFER_SIMULTANEOUS_USE_KHR, 0
+            };
+            command_buffer =
+                clCreateCommandBufferKHR(1, &queue, properties, &error);
+        }
+        else
+        {
+            command_buffer =
+                clCreateCommandBufferKHR(1, &queue, nullptr, &error);
+        }
+        test_error(error, "clCreateCommandBufferKHR failed");
+
+        return CL_SUCCESS;
+    }
+
+    // Test body returning an OpenCL error code
+    virtual cl_int Run() = 0;
+
+
+protected:
+    size_t data_size() const { return num_elements * sizeof(cl_int); }
+
+    cl_context context;
+    cl_command_queue queue;
+    clCommandBufferWrapper command_buffer;
+    clProgramWrapper program;
+    clKernelWrapper kernel;
+    clMemWrapper in_mem, out_mem;
+    size_t num_elements;
+
+    // Device support query results
+    bool simultaneous_use;
+    bool out_of_order_support;
+};
+
+// Test enqueuing a command-buffer containing a single NDRange command once
+struct BasicEnqueueTest : public BasicCommandBufferTest
+{
+    using BasicCommandBufferTest::BasicCommandBufferTest;
+
+    cl_int Run() override
+    {
+        cl_int error = clCommandNDRangeKernelKHR(
+            command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements,
+            nullptr, 0, nullptr, nullptr, nullptr);
+        test_error(error, "clCommandNDRangeKernelKHR failed");
+
+        error = clFinalizeCommandBufferKHR(command_buffer);
+        test_error(error, "clFinalizeCommandBufferKHR failed");
+
+        const cl_int pattern = 42;
+        error = clEnqueueFillBuffer(queue, in_mem, &pattern, sizeof(cl_int), 0,
+                                    data_size(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueFillBuffer failed");
+
+        error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
+                                          nullptr, nullptr);
+        test_error(error, "clEnqueueCommandBufferKHR failed");
+
+        std::vector<cl_int> output_data(num_elements);
+        error = clEnqueueReadBuffer(queue, out_mem, CL_TRUE, 0, data_size(),
+                                    output_data.data(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueReadBuffer failed");
+
+        for (size_t i = 0; i < num_elements; i++)
+        {
+            CHECK_VERIFICATION_ERROR(pattern, output_data[i], i);
+        }
+
+        return CL_SUCCESS;
+    }
+};
+
+// Test enqueuing a command-buffer containing multiple command, including
+// operations other than NDRange kernel execution.
+struct MixedCommandsTest : public BasicCommandBufferTest
+{
+    using BasicCommandBufferTest::BasicCommandBufferTest;
+
+    cl_int Run() override
+    {
+        cl_int error;
+        const size_t iterations = 4;
+        clMemWrapper result_mem =
+            clCreateBuffer(context, CL_MEM_READ_WRITE,
+                           sizeof(cl_int) * iterations, nullptr, &error);
+        test_error(error, "clCreateBuffer failed");
+
+        const cl_int pattern_base = 42;
+        for (size_t i = 0; i < iterations; i++)
+        {
+            const cl_int pattern = pattern_base + i;
+            cl_int error = clCommandFillBufferKHR(
+                command_buffer, nullptr, in_mem, &pattern, sizeof(cl_int), 0,
+                data_size(), 0, nullptr, nullptr, nullptr);
+            test_error(error, "clCommandFillBufferKHR failed");
+
+            error = clCommandNDRangeKernelKHR(
+                command_buffer, nullptr, nullptr, kernel, 1, nullptr,
+                &num_elements, nullptr, 0, nullptr, nullptr, nullptr);
+            test_error(error, "clCommandNDRangeKernelKHR failed");
+
+            const size_t result_offset = i * sizeof(cl_int);
+            error = clCommandCopyBufferKHR(
+                command_buffer, nullptr, out_mem, result_mem, 0, result_offset,
+                sizeof(cl_int), 0, nullptr, nullptr, nullptr);
+            test_error(error, "clCommandCopyBufferKHR failed");
+        }
+
+        error = clFinalizeCommandBufferKHR(command_buffer);
+        test_error(error, "clFinalizeCommandBufferKHR failed");
+
+        error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
+                                          nullptr, nullptr);
+        test_error(error, "clEnqueueCommandBufferKHR failed");
+
+        std::vector<cl_int> result_data(num_elements);
+        error = clEnqueueReadBuffer(queue, result_mem, CL_TRUE, 0,
+                                    iterations * sizeof(cl_int),
+                                    result_data.data(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueReadBuffer failed");
+
+        for (size_t i = 0; i < iterations; i++)
+        {
+            const cl_int ref = pattern_base + i;
+            CHECK_VERIFICATION_ERROR(ref, result_data[i], i);
+        }
+
+        return CL_SUCCESS;
+    }
+};
+
+// Test enqueueing a command-buffer blocked on a user-event
+struct UserEventTest : public BasicCommandBufferTest
+{
+    using BasicCommandBufferTest::BasicCommandBufferTest;
+
+    cl_int Run() override
+    {
+        cl_int error = clCommandNDRangeKernelKHR(
+            command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements,
+            nullptr, 0, nullptr, nullptr, nullptr);
+        test_error(error, "clCommandNDRangeKernelKHR failed");
+
+        error = clFinalizeCommandBufferKHR(command_buffer);
+        test_error(error, "clFinalizeCommandBufferKHR failed");
+
+        clEventWrapper user_event = clCreateUserEvent(context, &error);
+        test_error(error, "clCreateUserEvent failed");
+
+        const cl_int pattern = 42;
+        error = clEnqueueFillBuffer(queue, in_mem, &pattern, sizeof(cl_int), 0,
+                                    data_size(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueFillBuffer failed");
+
+        error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 1,
+                                          &user_event, nullptr);
+        test_error(error, "clEnqueueCommandBufferKHR failed");
+
+        std::vector<cl_int> output_data(num_elements);
+        error = clEnqueueReadBuffer(queue, out_mem, CL_FALSE, 0, data_size(),
+                                    output_data.data(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueReadBuffer failed");
+
+        error = clSetUserEventStatus(user_event, CL_COMPLETE);
+        test_error(error, "clSetUserEventStatus failed");
+
+        error = clFinish(queue);
+        test_error(error, "clFinish failed");
+
+        for (size_t i = 0; i < num_elements; i++)
+        {
+            CHECK_VERIFICATION_ERROR(pattern, output_data[i], i);
+        }
+
+        return CL_SUCCESS;
+    }
+};
+
+// Test flushing the command-queue between command-buffer enqueues
+struct ExplicitFlushTest : public BasicCommandBufferTest
+{
+    using BasicCommandBufferTest::BasicCommandBufferTest;
+
+    cl_int Run() override
+    {
+        cl_int error = clCommandNDRangeKernelKHR(
+            command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements,
+            nullptr, 0, nullptr, nullptr, nullptr);
+        test_error(error, "clCommandNDRangeKernelKHR failed");
+
+        error = clFinalizeCommandBufferKHR(command_buffer);
+        test_error(error, "clFinalizeCommandBufferKHR failed");
+
+        const cl_int pattern_A = 42;
+        error = clEnqueueFillBuffer(queue, in_mem, &pattern_A, sizeof(cl_int),
+                                    0, data_size(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueFillBuffer failed");
+
+        error = clFlush(queue);
+        test_error(error, "clFlush failed");
+
+        error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
+                                          nullptr, nullptr);
+        test_error(error, "clEnqueueCommandBufferKHR failed");
+
+        std::vector<cl_int> output_data_A(num_elements);
+        error = clEnqueueReadBuffer(queue, out_mem, CL_FALSE, 0, data_size(),
+                                    output_data_A.data(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueReadBuffer failed");
+
+        const cl_int pattern_B = 0xA;
+        error = clEnqueueFillBuffer(queue, in_mem, &pattern_B, sizeof(cl_int),
+                                    0, data_size(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueFillBuffer failed");
+
+        error = clFlush(queue);
+        test_error(error, "clFlush failed");
+
+        error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
+                                          nullptr, nullptr);
+        test_error(error, "clEnqueueCommandBufferKHR failed");
+
+        error = clFlush(queue);
+        test_error(error, "clFlush failed");
+
+        std::vector<cl_int> output_data_B(num_elements);
+        error = clEnqueueReadBuffer(queue, out_mem, CL_FALSE, 0, data_size(),
+                                    output_data_B.data(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueReadBuffer failed");
+
+        error = clFinish(queue);
+        test_error(error, "clFinish failed");
+
+        for (size_t i = 0; i < num_elements; i++)
+        {
+            CHECK_VERIFICATION_ERROR(pattern_A, output_data_A[i], i);
+
+            CHECK_VERIFICATION_ERROR(pattern_B, output_data_B[i], i);
+        }
+        return CL_SUCCESS;
+    }
+
+    bool Skip() override
+    {
+        return !simultaneous_use || BasicCommandBufferTest::Skip();
+    }
+};
+
+// Test enqueueing a command-buffer twice separated by another enqueue operation
+struct InterleavedEnqueueTest : public BasicCommandBufferTest
+{
+    using BasicCommandBufferTest::BasicCommandBufferTest;
+
+    cl_int Run() override
+    {
+        cl_int error = clCommandNDRangeKernelKHR(
+            command_buffer, nullptr, nullptr, kernel, 1, nullptr, &num_elements,
+            nullptr, 0, nullptr, nullptr, nullptr);
+        test_error(error, "clCommandNDRangeKernelKHR failed");
+
+        error = clFinalizeCommandBufferKHR(command_buffer);
+        test_error(error, "clFinalizeCommandBufferKHR failed");
+
+        cl_int pattern = 42;
+        error = clEnqueueFillBuffer(queue, in_mem, &pattern, sizeof(cl_int), 0,
+                                    data_size(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueFillBuffer failed");
+
+        error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
+                                          nullptr, nullptr);
+        test_error(error, "clEnqueueCommandBufferKHR failed");
+
+        pattern = 0xABCD;
+        error = clEnqueueFillBuffer(queue, in_mem, &pattern, sizeof(cl_int), 0,
+                                    data_size(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueFillBuffer failed");
+
+        error = clEnqueueCommandBufferKHR(0, nullptr, command_buffer, 0,
+                                          nullptr, nullptr);
+        test_error(error, "clEnqueueCommandBufferKHR failed");
+
+        error = clEnqueueCopyBuffer(queue, in_mem, out_mem, 0, 0, data_size(),
+                                    0, nullptr, nullptr);
+        test_error(error, "clEnqueueCopyBuffer failed");
+
+        std::vector<cl_int> output_data(num_elements);
+        error = clEnqueueReadBuffer(queue, out_mem, CL_TRUE, 0, data_size(),
+                                    output_data.data(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueReadBuffer failed");
+
+        for (size_t i = 0; i < num_elements; i++)
+        {
+            CHECK_VERIFICATION_ERROR(pattern, output_data[i], i);
+        }
+
+        return CL_SUCCESS;
+    }
+
+    bool Skip() override
+    {
+        return !simultaneous_use || BasicCommandBufferTest::Skip();
+    }
+};
+
+// Test sync-points with an out-of-order command-buffer
+struct OutOfOrderTest : public BasicCommandBufferTest
+{
+    using BasicCommandBufferTest::BasicCommandBufferTest;
+    OutOfOrderTest(cl_device_id device, cl_context context,
+                   cl_command_queue queue)
+        : BasicCommandBufferTest(device, context, queue),
+          out_of_order_command_buffer(this), out_of_order_queue(nullptr),
+          event(nullptr)
+    {}
+
+    cl_int Run() override
+    {
+        cl_sync_point_khr sync_points[2];
+
+        const cl_int pattern = 42;
+        cl_int error =
+            clCommandFillBufferKHR(out_of_order_command_buffer, nullptr, in_mem,
+                                   &pattern, sizeof(cl_int), 0, data_size(), 0,
+                                   nullptr, &sync_points[0], nullptr);
+        test_error(error, "clCommandFillBufferKHR failed");
+
+        const cl_int overwritten_pattern = 0xACDC;
+        error = clCommandFillBufferKHR(out_of_order_command_buffer, nullptr,
+                                       out_mem, &overwritten_pattern,
+                                       sizeof(cl_int), 0, data_size(), 0,
+                                       nullptr, &sync_points[1], nullptr);
+        test_error(error, "clCommandFillBufferKHR failed");
+
+        error = clCommandNDRangeKernelKHR(
+            out_of_order_command_buffer, nullptr, nullptr, kernel, 1, nullptr,
+            &num_elements, nullptr, 2, sync_points, nullptr, nullptr);
+        test_error(error, "clCommandNDRangeKernelKHR failed");
+
+        error = clFinalizeCommandBufferKHR(out_of_order_command_buffer);
+        test_error(error, "clFinalizeCommandBufferKHR failed");
+
+        error = clEnqueueCommandBufferKHR(
+            0, nullptr, out_of_order_command_buffer, 0, nullptr, &event);
+        test_error(error, "clEnqueueCommandBufferKHR failed");
+
+        std::vector<cl_int> output_data(num_elements);
+        error = clEnqueueReadBuffer(out_of_order_queue, out_mem, CL_TRUE, 0,
+                                    data_size(), output_data.data(), 1, &event,
+                                    nullptr);
+        test_error(error, "clEnqueueReadBuffer failed");
+
+        for (size_t i = 0; i < num_elements; i++)
+        {
+            CHECK_VERIFICATION_ERROR(pattern, output_data[i], i);
+        }
+
+        return CL_SUCCESS;
+    }
+
+    cl_int SetUp(int elements) override
+    {
+        cl_int error = BasicCommandBufferTest::SetUp(elements);
+        test_error(error, "BasicCommandBufferTest::SetUp failed");
+
+        if (!out_of_order_support)
+        {
+            // Test will skip as device doesn't support out-of-order
+            // command-buffers
+            return CL_SUCCESS;
+        }
+
+        out_of_order_queue = clCreateCommandQueue(
+            context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &error);
+        test_error(error, "Unable to create command queue to test with");
+
+        out_of_order_command_buffer =
+            clCreateCommandBufferKHR(1, &out_of_order_queue, nullptr, &error);
+        test_error(error, "clCreateCommandBufferKHR failed");
+
+        return CL_SUCCESS;
+    }
+
+    bool Skip() override
+    {
+        return !out_of_order_support || BasicCommandBufferTest::Skip();
+    }
+
+    clCommandQueueWrapper out_of_order_queue;
+    clCommandBufferWrapper out_of_order_command_buffer;
+    clEventWrapper event;
+};
+
+#undef CHECK_VERIFICATION_ERROR
+
+template <class T>
+int MakeAndRunTest(cl_device_id device, cl_context context,
+                   cl_command_queue queue, int num_elements)
+{
+    CHECK_COMMAND_BUFFER_EXTENSION_AVAILABLE(device);
+
+    auto test_fixture = T(device, context, queue);
+    cl_int error = test_fixture.SetUp(num_elements);
+    test_error_ret(error, "Error in test initialization", TEST_FAIL);
+
+    if (test_fixture.Skip())
+    {
+        return TEST_SKIPPED_ITSELF;
+    }
+
+    error = test_fixture.Run();
+    test_error_ret(error, "Test Failed", TEST_FAIL);
+
+    return TEST_PASS;
+}
+} // anonymous namespace
+
+int test_single_ndrange(cl_device_id device, cl_context context,
+                        cl_command_queue queue, int num_elements)
+{
+    return MakeAndRunTest<BasicEnqueueTest>(device, context, queue,
+                                            num_elements);
+}
+
+int test_interleaved_enqueue(cl_device_id device, cl_context context,
+                             cl_command_queue queue, int num_elements)
+{
+    return MakeAndRunTest<InterleavedEnqueueTest>(device, context, queue,
+                                                  num_elements);
+}
+
+int test_mixed_commands(cl_device_id device, cl_context context,
+                        cl_command_queue queue, int num_elements)
+{
+    return MakeAndRunTest<MixedCommandsTest>(device, context, queue,
+                                             num_elements);
+}
+
+int test_explicit_flush(cl_device_id device, cl_context context,
+                        cl_command_queue queue, int num_elements)
+{
+    return MakeAndRunTest<ExplicitFlushTest>(device, context, queue,
+                                             num_elements);
+}
+
+int test_user_events(cl_device_id device, cl_context context,
+                     cl_command_queue queue, int num_elements)
+{
+    return MakeAndRunTest<UserEventTest>(device, context, queue, num_elements);
+}
+
+int test_out_of_order(cl_device_id device, cl_context context,
+                      cl_command_queue queue, int num_elements)
+{
+    return MakeAndRunTest<OutOfOrderTest>(device, context, queue, num_elements);
+}

test_conformance/extensions/cl_khr_command_buffer/command_buffer_test_base.h

@@ -0,0 +1,177 @@
+//
+// Copyright (c) 2022 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef _CL_KHR_COMMAND_BUFFER_TEST_BASE_H
+#define _CL_KHR_COMMAND_BUFFER_TEST_BASE_H
+
+#include <CL/cl_ext.h>
+#include "harness/deviceInfo.h"
+#include "harness/testHarness.h"
+
+
+// Base class for setting function pointers to new extension entry points
+struct CommandBufferTestBase
+{
+    CommandBufferTestBase(cl_device_id device): device(device) {}
+
+    cl_int init_extension_functions()
+    {
+        cl_platform_id platform;
+        cl_int error =
+            clGetDeviceInfo(device, CL_DEVICE_PLATFORM, sizeof(cl_platform_id),
+                            &platform, nullptr);
+        test_error(error, "clGetDeviceInfo for CL_DEVICE_PLATFORM failed");
+
+        // If it is supported get the addresses of all the APIs here.
+#define GET_EXTENSION_ADDRESS(FUNC)                                            \
+    FUNC = reinterpret_cast<FUNC##_fn>(                                        \
+        clGetExtensionFunctionAddressForPlatform(platform, #FUNC));            \
+    if (FUNC == nullptr)                                                       \
+    {                                                                          \
+        log_error("ERROR: clGetExtensionFunctionAddressForPlatform failed"     \
+                  " with " #FUNC "\n");                                        \
+        return TEST_FAIL;                                                      \
+    }
+
+        GET_EXTENSION_ADDRESS(clCreateCommandBufferKHR);
+        GET_EXTENSION_ADDRESS(clReleaseCommandBufferKHR);
+        GET_EXTENSION_ADDRESS(clRetainCommandBufferKHR);
+        GET_EXTENSION_ADDRESS(clFinalizeCommandBufferKHR);
+        GET_EXTENSION_ADDRESS(clEnqueueCommandBufferKHR);
+        GET_EXTENSION_ADDRESS(clCommandBarrierWithWaitListKHR);
+        GET_EXTENSION_ADDRESS(clCommandCopyBufferKHR);
+        GET_EXTENSION_ADDRESS(clCommandCopyBufferRectKHR);
+        GET_EXTENSION_ADDRESS(clCommandCopyBufferToImageKHR);
+        GET_EXTENSION_ADDRESS(clCommandCopyImageKHR);
+        GET_EXTENSION_ADDRESS(clCommandCopyImageToBufferKHR);
+        GET_EXTENSION_ADDRESS(clCommandFillBufferKHR);
+        GET_EXTENSION_ADDRESS(clCommandFillImageKHR);
+        GET_EXTENSION_ADDRESS(clCommandNDRangeKernelKHR);
+        GET_EXTENSION_ADDRESS(clGetCommandBufferInfoKHR);
+#undef GET_EXTENSION_ADDRESS
+        return CL_SUCCESS;
+    }
+
+    clCreateCommandBufferKHR_fn clCreateCommandBufferKHR = nullptr;
+    clReleaseCommandBufferKHR_fn clReleaseCommandBufferKHR = nullptr;
+    clRetainCommandBufferKHR_fn clRetainCommandBufferKHR = nullptr;
+    clFinalizeCommandBufferKHR_fn clFinalizeCommandBufferKHR = nullptr;
+    clEnqueueCommandBufferKHR_fn clEnqueueCommandBufferKHR = nullptr;
+    clCommandBarrierWithWaitListKHR_fn clCommandBarrierWithWaitListKHR =
+        nullptr;
+    clCommandCopyBufferKHR_fn clCommandCopyBufferKHR = nullptr;
+    clCommandCopyBufferRectKHR_fn clCommandCopyBufferRectKHR = nullptr;
+    clCommandCopyBufferToImageKHR_fn clCommandCopyBufferToImageKHR = nullptr;
+    clCommandCopyImageKHR_fn clCommandCopyImageKHR = nullptr;
+    clCommandCopyImageToBufferKHR_fn clCommandCopyImageToBufferKHR = nullptr;
+    clCommandFillBufferKHR_fn clCommandFillBufferKHR = nullptr;
+    clCommandFillImageKHR_fn clCommandFillImageKHR = nullptr;
+    clCommandNDRangeKernelKHR_fn clCommandNDRangeKernelKHR = nullptr;
+    clGetCommandBufferInfoKHR_fn clGetCommandBufferInfoKHR = nullptr;
+
+    cl_device_id device = nullptr;
+};
+
+// Wrapper class based off generic typeWrappers.h wrappers. However, because
+// the release/retain functions are queried at runtime from the platform,
+// rather than known at compile time we cannot link the instantiated template.
+// Instead, pass an instance of `CommandBufferTestBase` on wrapper construction
+// to access the release/retain functions.
+class clCommandBufferWrapper {
+    cl_command_buffer_khr object = nullptr;
+
+    void retain()
+    {
+        if (!object) return;
+
+        auto err = base->clRetainCommandBufferKHR(object);
+        if (err != CL_SUCCESS)
+        {
+            print_error(err, "clRetainCommandBufferKHR() failed");
+            std::abort();
+        }
+    }
+
+    void release()
+    {
+        if (!object) return;
+
+        auto err = base->clReleaseCommandBufferKHR(object);
+        if (err != CL_SUCCESS)
+        {
+            print_error(err, "clReleaseCommandBufferKHR() failed");
+            std::abort();
+        }
+    }
+
+    // Used to access release/retain functions
+    CommandBufferTestBase *base;
+
+public:
+    // We always want to have base available to dereference
+    clCommandBufferWrapper() = delete;
+
+    clCommandBufferWrapper(CommandBufferTestBase *base): base(base) {}
+
+    // On assignment, assume the object has a refcount of one.
+    clCommandBufferWrapper &operator=(cl_command_buffer_khr rhs)
+    {
+        reset(rhs);
+        return *this;
+    }
+
+    // Copy semantics, increase retain count.
+    clCommandBufferWrapper(clCommandBufferWrapper const &w) { *this = w; }
+    clCommandBufferWrapper &operator=(clCommandBufferWrapper const &w)
+    {
+        reset(w.object);
+        retain();
+        return *this;
+    }
+
+    // Move semantics, directly take ownership.
+    clCommandBufferWrapper(clCommandBufferWrapper &&w) { *this = std::move(w); }
+    clCommandBufferWrapper &operator=(clCommandBufferWrapper &&w)
+    {
+        reset(w.object);
+        w.object = nullptr;
+        return *this;
+    }
+
+    ~clCommandBufferWrapper() { reset(); }
+
+    // Release the existing object, if any, and own the new one, if any.
+    void reset(cl_command_buffer_khr new_object = nullptr)
+    {
+        release();
+        object = new_object;
+    }
+
+    operator cl_command_buffer_khr() const { return object; }
+};
+
+#define CHECK_COMMAND_BUFFER_EXTENSION_AVAILABLE(device)                       \
+    {                                                                          \
+        if (!is_extension_available(device, "cl_khr_command_buffer"))          \
+        {                                                                      \
+            log_info(                                                          \
+                "Device does not support 'cl_khr_command_buffer'. Skipping "   \
+                "the test.\n");                                                \
+            return TEST_SKIPPED_ITSELF;                                        \
+        }                                                                      \
+    }
+
+
+#endif // _CL_KHR_COMMAND_BUFFER_TEST_BASE_H

test_conformance/extensions/cl_khr_command_buffer/main.cpp

@@ -0,0 +1,35 @@
+// Copyright (c) 2022 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+#include "procs.h"
+#include "harness/testHarness.h"
+
+test_definition test_list[] = {
+    ADD_TEST(single_ndrange), ADD_TEST(interleaved_enqueue),
+    ADD_TEST(mixed_commands), ADD_TEST(explicit_flush),
+    ADD_TEST(user_events),    ADD_TEST(out_of_order)
+};
+
+
+int main(int argc, const char *argv[])
+{
+    // A device may report the required properties of a queue that
+    // is compatible with command-buffers via the query
+    // CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR. We account
+    // for this in the tests themselves, rather than here, where we have a
+    // device to query.
+    const cl_command_queue_properties queue_properties = 0;
+    return runTestHarnessWithCheck(argc, argv, ARRAY_SIZE(test_list), test_list,
+                                   false, queue_properties, nullptr);
+}

test_conformance/extensions/cl_khr_command_buffer/procs.h

@@ -0,0 +1,35 @@
+//
+// Copyright (c) 2022 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+#ifndef _CL_KHR_COMMAND_BUFFER_PROCS_H
+#define _CL_KHR_COMMAND_BUFFER_PROCS_H
+
+#include <CL/cl.h>
+
+// Basic command-buffer tests
+extern int test_single_ndrange(cl_device_id device, cl_context context,
+                               cl_command_queue queue, int num_elements);
+extern int test_interleaved_enqueue(cl_device_id device, cl_context context,
+                                    cl_command_queue queue, int num_elements);
+extern int test_mixed_commands(cl_device_id device, cl_context context,
+                               cl_command_queue queue, int num_elements);
+extern int test_explicit_flush(cl_device_id device, cl_context context,
+                               cl_command_queue queue, int num_elements);
+extern int test_user_events(cl_device_id device, cl_context context,
+                            cl_command_queue queue, int num_elements);
+extern int test_out_of_order(cl_device_id device, cl_context context,
+                             cl_command_queue queue, int num_elements);
+
+#endif /*_CL_KHR_COMMAND_BUFFER_PROCS_H*/