Command buffer queue substitution (#1584)

* Added queue substitution CTS test based on issue: https://github.com/KhronosGroup/OpenCL-CTS/issues/1369

* Added cosmetic correction to properties_queue_substitution test, work in progress:

https://github.com/KhronosGroup/OpenCL-CTS/issues/1369

* Added support for simultaneous queue substitution test (1.3.3), more cosmetic corrections added:

https://github.com/KhronosGroup/OpenCL-CTS/issues/1369

* Added initial rearrangement of source files due to more test cases to be implemented.

* Corrected test with queue properties out-of-order:

https://github.com/KhronosGroup/OpenCL-CTS/issues/1369

* Added alternative temporary path for simultenaous pass (USE_COMMAND_BUF_KENEL_ARG):

https://github.com/KhronosGroup/OpenCL-CTS/issues/1369

* applied git clang-format "origin/main"

* Added corrections for out-of-order property scenario, couple cosmetic fixes:

https://github.com/KhronosGroup/OpenCL-CTS/issues/1369

* Replaced cl_command_queue with clCommandQueueWrapper for additional command queue with properties:

https://github.com/KhronosGroup/OpenCL-CTS/issues/1369

* Removed unnecessary flag, additional cleanup:

https://github.com/KhronosGroup/OpenCL-CTS/issues/1369

* Added narrowing correction due to windows build fail:

https://github.com/KhronosGroup/OpenCL-CTS/issues/1369

* Test simiplified in order to use only substitute queue on both simultaneous buffers:

https://github.com/KhronosGroup/OpenCL-CTS/issues/1369

* Added limitation to use only CL_QUEUE_PROFILING_ENABLE property, replaced cl_command_queue with related wrapper (#1369, p.1.4)

* Corrections related to order of operations due to code review (issue #1369, p.1.3)

* Cosmetic fix for prev commit

* Cosmetic fix for prev commit #2

test_conformance/extensions/cl_khr_command_buffer/CMakeLists.txt

@@ -3,6 +3,7 @@ set(MODULE_NAME CL_KHR_COMMAND_BUFFER)
 set(${MODULE_NAME}_SOURCES
     main.cpp
     basic_command_buffer.cpp
+    command_buffer_queue_substitution.cpp
 )
 
 include(../../CMakeCommon.txt)

test_conformance/extensions/cl_khr_command_buffer/basic_command_buffer.cpp

@@ -13,158 +13,148 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 //
-#include "command_buffer_test_base.h"
+#include "basic_command_buffer.h"
 #include "procs.h"
-#include "harness/typeWrappers.h"
 
 #include <algorithm>
 #include <cstring>
 #include <vector>
 
-#define CHECK_VERIFICATION_ERROR(reference, result, index)                     \
+
-    {                                                                          \
+BasicCommandBufferTest::BasicCommandBufferTest(cl_device_id device,
-        if (reference != result)                                               \
+                                               cl_context context,
-        {                                                                      \
+                                               cl_command_queue queue)
-            log_error("Expected %d was %d at index %u\n", reference, result,   \
+    : CommandBufferTestBase(device), context(context), queue(nullptr),
-                      index);                                                  \
+      num_elements(0), command_buffer(this), simultaneous_use_support(false),
-            return TEST_FAIL;                                                  \
+      out_of_order_support(false),
-        }                                                                      \
+      // try to use simultaneous path by default
+      simultaneous_use_requested(true),
+      // due to simultaneous cases extend buffer size
+      buffer_size_multiplier(1)
+
+{
+    cl_int error = clRetainCommandQueue(queue);
+    if (error != CL_SUCCESS)
+    {
+        throw std::runtime_error("clRetainCommandQueue failed\n");
     }
+    this->queue = queue;
+}
+
+bool BasicCommandBufferTest::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");
+
+
+    // 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_support = simultaneous_use_requested
+        && (capabilities & CL_COMMAND_BUFFER_CAPABILITY_SIMULTANEOUS_USE_KHR)
+            != 0;
+    out_of_order_support =
+        capabilities & CL_COMMAND_BUFFER_CAPABILITY_OUT_OF_ORDER_KHR;
+
+    // Skip if queue properties don't contain those required
+    return required_properties != (required_properties & queue_properties);
+}
+
+cl_int BasicCommandBufferTest::SetUp(int elements)
+{
+    cl_int error = init_extension_functions();
+    if (error != CL_SUCCESS)
+    {
+        return error;
+    }
+
+    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, __global int* offset) {
+        size_t id = get_global_id(0);
+        int ind = offset[0] + id;
+        out[ind] = in[ind];
+    })";
+
+    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 * buffer_size_multiplier,
+                       nullptr, &error);
+    test_error(error, "clCreateBuffer failed");
+
+    out_mem =
+        clCreateBuffer(context, CL_MEM_WRITE_ONLY,
+                       sizeof(cl_int) * num_elements * buffer_size_multiplier,
+                       nullptr, &error);
+    test_error(error, "clCreateBuffer failed");
+
+    cl_int offset = 0;
+    off_mem = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
+                             sizeof(cl_int), &offset, &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");
+
+    error = clSetKernelArg(kernel, 2, sizeof(off_mem), &off_mem);
+    test_error(error, "clSetKernelArg failed");
+
+    if (simultaneous_use_support)
+    {
+        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;
+}
 
 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), num_elements(0), simultaneous_use(false),
-          out_of_order_support(false)
-    {}
-
-    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
 {
@@ -375,7 +365,7 @@ struct ExplicitFlushTest : public BasicCommandBufferTest
 
     bool Skip() override
     {
-        return !simultaneous_use || BasicCommandBufferTest::Skip();
+        return BasicCommandBufferTest::Skip() || !simultaneous_use_support;
     }
 };
 
@@ -431,7 +421,7 @@ struct InterleavedEnqueueTest : public BasicCommandBufferTest
 
     bool Skip() override
     {
-        return !simultaneous_use || BasicCommandBufferTest::Skip();
+        return BasicCommandBufferTest::Skip() || !simultaneous_use_support;
     }
 };
 
@@ -495,13 +485,6 @@ struct OutOfOrderTest : public BasicCommandBufferTest
         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");
@@ -515,7 +498,7 @@ struct OutOfOrderTest : public BasicCommandBufferTest
 
     bool Skip() override
     {
-        return !out_of_order_support || BasicCommandBufferTest::Skip();
+        return BasicCommandBufferTest::Skip() || !out_of_order_support;
     }
 
     clCommandQueueWrapper out_of_order_queue;
@@ -523,28 +506,6 @@ struct OutOfOrderTest : public BasicCommandBufferTest
     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,

test_conformance/extensions/cl_khr_command_buffer/basic_command_buffer.h

@@ -0,0 +1,100 @@
+//
+// 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_BASIC_COMMAND_BUFFER_H
+#define _CL_KHR_BASIC_COMMAND_BUFFER_H
+
+#include "command_buffer_test_base.h"
+#include "harness/typeWrappers.h"
+
+#define ADD_PROP(prop)                                                         \
+    {                                                                          \
+        prop, #prop                                                            \
+    }
+
+#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;                                                  \
+        }                                                                      \
+    }
+
+// 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);
+
+    virtual bool Skip();
+    virtual cl_int SetUp(int elements);
+
+    // 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;
+    clCommandQueueWrapper queue;
+    clCommandBufferWrapper command_buffer;
+    clProgramWrapper program;
+    clKernelWrapper kernel;
+    clMemWrapper in_mem, out_mem, off_mem;
+    size_t num_elements;
+
+    // Device support query results
+    bool simultaneous_use_support;
+    bool out_of_order_support;
+
+    // user request for simultaneous use
+    bool simultaneous_use_requested;
+    unsigned buffer_size_multiplier;
+};
+
+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);
+
+    try
+    {
+        auto test_fixture = T(device, context, queue);
+
+        if (test_fixture.Skip())
+        {
+            return TEST_SKIPPED_ITSELF;
+        }
+
+        cl_int error = test_fixture.SetUp(num_elements);
+        test_error_ret(error, "Error in test initialization", TEST_FAIL);
+
+        error = test_fixture.Run();
+        test_error_ret(error, "Test Failed", TEST_FAIL);
+    } catch (const std::runtime_error &e)
+    {
+        log_error("%s", e.what());
+        return TEST_FAIL;
+    }
+
+    return TEST_PASS;
+}
+
+#endif // _CL_KHR_BASIC_COMMAND_BUFFER_H

test_conformance/extensions/cl_khr_command_buffer/command_buffer_queue_substitution.cpp

@@ -0,0 +1,278 @@
+//
+// 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 "basic_command_buffer.h"
+#include "procs.h"
+
+#include <vector>
+
+namespace {
+
+////////////////////////////////////////////////////////////////////////////////
+// Command-queue substitution tests which handles below cases:
+// -substitution on queue without properties
+// -substitution on queue with properties
+// -simultaneous use queue substitution
+
+template <bool prop_use, bool simul_use>
+struct SubstituteQueueTest : public BasicCommandBufferTest
+{
+    SubstituteQueueTest(cl_device_id device, cl_context context,
+                        cl_command_queue queue)
+        : BasicCommandBufferTest(device, context, queue),
+          properties_use_requested(prop_use), user_event(nullptr)
+    {
+        simultaneous_use_requested = simul_use;
+        if (simul_use) buffer_size_multiplier = 2;
+    }
+
+    //--------------------------------------------------------------------------
+    bool Skip() override
+    {
+        if (properties_use_requested)
+        {
+            Version version = get_device_cl_version(device);
+            const cl_device_info host_queue_query = version >= Version(2, 0)
+                ? CL_DEVICE_QUEUE_ON_HOST_PROPERTIES
+                : CL_DEVICE_QUEUE_PROPERTIES;
+
+            cl_queue_properties host_queue_props = 0;
+            int error = clGetDeviceInfo(device, host_queue_query,
+                                        sizeof(host_queue_props),
+                                        &host_queue_props, NULL);
+            test_error(error, "clGetDeviceInfo failed");
+
+            if ((host_queue_props & CL_QUEUE_PROFILING_ENABLE) == 0)
+                return true;
+        }
+
+        return BasicCommandBufferTest::Skip()
+            || (simultaneous_use_requested && !simultaneous_use_support);
+    }
+
+    //--------------------------------------------------------------------------
+    cl_int SetUp(int elements) override
+    {
+        // By default command queue is created without properties,
+        // if test requires queue with properties default queue must be
+        // replaced.
+        if (properties_use_requested)
+        {
+            // due to the skip condition
+            cl_int error = CL_SUCCESS;
+            queue = clCreateCommandQueue(context, device,
+                                         CL_QUEUE_PROFILING_ENABLE, &error);
+            test_error(
+                error,
+                "clCreateCommandQueue with CL_QUEUE_PROFILING_ENABLE failed");
+        }
+
+        return BasicCommandBufferTest::SetUp(elements);
+    }
+
+    //--------------------------------------------------------------------------
+    cl_int Run() override
+    {
+        // record command buffer with primary queue
+        cl_int error = RecordCommandBuffer();
+        test_error(error, "RecordCommandBuffer failed");
+
+        // create substitute queue
+        clCommandQueueWrapper new_queue;
+        if (properties_use_requested)
+        {
+            new_queue = clCreateCommandQueue(context, device,
+                                             CL_QUEUE_PROFILING_ENABLE, &error);
+            test_error(
+                error,
+                "clCreateCommandQueue with CL_QUEUE_PROFILING_ENABLE failed");
+        }
+        else
+        {
+            const cl_command_queue_properties queue_properties = 0;
+            new_queue =
+                clCreateCommandQueue(context, device, queue_properties, &error);
+            test_error(error, "clCreateCommandQueue failed");
+        }
+
+        if (simultaneous_use_support)
+        {
+            // enque simultaneous command-buffers with substitute queue
+            error = RunSimultaneous(new_queue);
+            test_error(error, "RunSimultaneous failed");
+        }
+        else
+        {
+            // enque single command-buffer with substitute queue
+            error = RunSingle(new_queue);
+            test_error(error, "RunSingle failed");
+        }
+
+        return CL_SUCCESS;
+    }
+
+    //--------------------------------------------------------------------------
+    cl_int RecordCommandBuffer()
+    {
+        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");
+        return CL_SUCCESS;
+    }
+
+    //--------------------------------------------------------------------------
+    cl_int RunSingle(const cl_command_queue& q)
+    {
+        cl_int error = CL_SUCCESS;
+        std::vector<cl_int> output_data(num_elements);
+
+        error = clEnqueueFillBuffer(q, in_mem, &pattern_pri, sizeof(cl_int), 0,
+                                    data_size(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueFillBuffer failed");
+
+        cl_command_queue queues[] = { q };
+        error = clEnqueueCommandBufferKHR(1, queues, command_buffer, 0, nullptr,
+                                          nullptr);
+        test_error(error, "clEnqueueCommandBufferKHR failed");
+
+        error = clEnqueueReadBuffer(q, out_mem, CL_TRUE, 0, data_size(),
+                                    output_data.data(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueReadBuffer failed");
+
+        error = clFinish(q);
+        test_error(error, "clFinish failed");
+
+        for (size_t i = 0; i < num_elements; i++)
+        {
+            CHECK_VERIFICATION_ERROR(pattern_pri, output_data[i], i);
+        }
+
+        return CL_SUCCESS;
+    }
+
+    //--------------------------------------------------------------------------
+    struct SimulPassData
+    {
+        cl_int pattern;
+        cl_int offset;
+        cl_command_queue queue;
+        std::vector<cl_int> output_buffer;
+    };
+
+    //--------------------------------------------------------------------------
+    cl_int EnqueueSimultaneousPass(SimulPassData& pd)
+    {
+        cl_int error = clEnqueueFillBuffer(
+            pd.queue, in_mem, &pd.pattern, sizeof(cl_int),
+            pd.offset * sizeof(cl_int), data_size(), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueFillBuffer failed");
+
+        error =
+            clEnqueueFillBuffer(pd.queue, off_mem, &pd.offset, sizeof(cl_int),
+                                0, sizeof(cl_int), 0, nullptr, nullptr);
+        test_error(error, "clEnqueueFillBuffer failed");
+
+        if (!user_event)
+        {
+            user_event = clCreateUserEvent(context, &error);
+            test_error(error, "clCreateUserEvent failed");
+        }
+
+        cl_command_queue queues[] = { pd.queue };
+        error = clEnqueueCommandBufferKHR(1, queues, command_buffer, 1,
+                                          &user_event, nullptr);
+        test_error(error, "clEnqueueCommandBufferKHR failed");
+
+        error = clEnqueueReadBuffer(
+            pd.queue, out_mem, CL_FALSE, pd.offset * sizeof(cl_int),
+            data_size(), pd.output_buffer.data(), 0, nullptr, nullptr);
+
+        test_error(error, "clEnqueueReadBuffer failed");
+
+        return CL_SUCCESS;
+    }
+
+    //--------------------------------------------------------------------------
+    cl_int RunSimultaneous(const cl_command_queue& q)
+    {
+        cl_int error = CL_SUCCESS;
+        cl_int offset = static_cast<cl_int>(num_elements);
+
+        std::vector<SimulPassData> simul_passes = {
+            { pattern_pri, 0, q, std::vector<cl_int>(num_elements) },
+            { pattern_sec, offset, q, std::vector<cl_int>(num_elements) }
+        };
+
+        for (auto&& pass : simul_passes)
+        {
+            error = EnqueueSimultaneousPass(pass);
+            test_error(error, "EnqueuePass failed");
+        }
+
+        error = clSetUserEventStatus(user_event, CL_COMPLETE);
+        test_error(error, "clSetUserEventStatus failed");
+
+        for (auto&& pass : simul_passes)
+        {
+            error = clFinish(pass.queue);
+            test_error(error, "clFinish failed");
+
+            auto& res_data = pass.output_buffer;
+
+            for (size_t i = 0; i < num_elements; i++)
+            {
+                CHECK_VERIFICATION_ERROR(pass.pattern, res_data[i], i);
+            }
+        }
+
+        return CL_SUCCESS;
+    }
+
+    //--------------------------------------------------------------------------
+    const cl_int pattern_pri = 0xB;
+    const cl_int pattern_sec = 0xC;
+
+    bool properties_use_requested;
+    clEventWrapper user_event;
+};
+
+} // anonymous namespace
+
+int test_queue_substitution(cl_device_id device, cl_context context,
+                            cl_command_queue queue, int num_elements)
+{
+    return MakeAndRunTest<SubstituteQueueTest<false, false>>(
+        device, context, queue, num_elements);
+}
+
+int test_properties_queue_substitution(cl_device_id device, cl_context context,
+                                       cl_command_queue queue, int num_elements)
+{
+    return MakeAndRunTest<SubstituteQueueTest<true, false>>(
+        device, context, queue, num_elements);
+}
+
+int test_simultaneous_queue_substitution(cl_device_id device,
+                                         cl_context context,
+                                         cl_command_queue queue,
+                                         int num_elements)
+{
+    return MakeAndRunTest<SubstituteQueueTest<false, true>>(
+        device, context, queue, num_elements);
+}

test_conformance/extensions/cl_khr_command_buffer/main.cpp

@@ -15,11 +15,15 @@
 #include "procs.h"
 #include "harness/testHarness.h"
 
-test_definition test_list[] = {
+test_definition test_list[] = { ADD_TEST(single_ndrange),
-    ADD_TEST(single_ndrange), ADD_TEST(interleaved_enqueue),
+                                ADD_TEST(interleaved_enqueue),
-    ADD_TEST(mixed_commands), ADD_TEST(explicit_flush),
+                                ADD_TEST(mixed_commands),
-    ADD_TEST(user_events),    ADD_TEST(out_of_order)
+                                ADD_TEST(explicit_flush),
-};
+                                ADD_TEST(user_events),
+                                ADD_TEST(out_of_order),
+                                ADD_TEST(queue_substitution),
+                                ADD_TEST(properties_queue_substitution),
+                                ADD_TEST(simultaneous_queue_substitution) };
 
 
 int main(int argc, const char *argv[])

test_conformance/extensions/cl_khr_command_buffer/procs.h

@@ -31,5 +31,16 @@ 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);
+extern int test_queue_substitution(cl_device_id device, cl_context context,
+                                   cl_command_queue queue, int num_elements);
+extern int test_properties_queue_substitution(cl_device_id device,
+                                              cl_context context,
+                                              cl_command_queue queue,
+                                              int num_elements);
+extern int test_simultaneous_queue_substitution(cl_device_id device,
+                                                cl_context context,
+                                                cl_command_queue queue,
+                                                int num_elements);
+
 
 #endif /*_CL_KHR_COMMAND_BUFFER_PROCS_H*/