Added semaphore test to cover out-of-order operations synchronized with barrier (#2068)

Fixes #1979 according to task description.

test_conformance/extensions/cl_khr_semaphore/CMakeLists.txt

@@ -7,6 +7,7 @@ set(${MODULE_NAME}_SOURCES
          test_semaphores_negative_release_retain.cpp
          test_semaphores_negative_getinfo.cpp
          test_semaphores_negative_create.cpp
+         test_semaphores_cross_queue.cpp
          test_semaphores_queries.cpp
          semaphore_base.h
 )

test_conformance/extensions/cl_khr_semaphore/main.cpp

@@ -37,6 +37,8 @@ test_definition test_list[] = {
     ADD_TEST_VERSION(semaphores_no_device_list_queries, Version(1, 2)),
     ADD_TEST_VERSION(semaphores_multi_device_context_queries, Version(1, 2)),
     ADD_TEST_VERSION(semaphores_import_export_fd, Version(1, 2)),
+    ADD_TEST_VERSION(semaphores_ooo_ops_single_queue, Version(1, 2)),
+    ADD_TEST_VERSION(semaphores_ooo_ops_cross_queue, Version(1, 2)),
     ADD_TEST_VERSION(semaphores_negative_create_invalid_context, Version(1, 2)),
     ADD_TEST_VERSION(semaphores_negative_create_invalid_property,
                      Version(1, 2)),

test_conformance/extensions/cl_khr_semaphore/procs.h

@@ -130,3 +130,11 @@ extern int test_semaphores_negative_retain(cl_device_id device,
                                            cl_context context,
                                            cl_command_queue queue,
                                            int num_elements);
+extern int test_semaphores_ooo_ops_single_queue(cl_device_id deviceID,
+                                                cl_context context,
+                                                cl_command_queue defaultQueue,
+                                                int num_elements);
+extern int test_semaphores_ooo_ops_cross_queue(cl_device_id deviceID,
+                                               cl_context context,
+                                               cl_command_queue defaultQueue,
+                                               int num_elements);

test_conformance/extensions/cl_khr_semaphore/semaphore_base.h

@@ -175,8 +175,9 @@ public:
 struct SemaphoreTestBase : public SemaphoreBase
 {
     SemaphoreTestBase(cl_device_id device, cl_context context,
-                      cl_command_queue queue)
+                      cl_command_queue queue, cl_int nelems)
-        : SemaphoreBase(device), context(context), semaphore(this)
+        : SemaphoreBase(device), context(context), semaphore(this),
+          num_elems(nelems)
     {
         cl_int error = init_extension_functions();
         if (error != CL_SUCCESS)
@@ -194,11 +195,12 @@ protected:
     cl_context context = nullptr;
     clCommandQueueWrapper queue = nullptr;
     clSemaphoreWrapper semaphore = nullptr;
+    cl_int num_elems = 0;
 };
 
 template <class T>
 int MakeAndRunTest(cl_device_id device, cl_context context,
-                   cl_command_queue queue)
+                   cl_command_queue queue, cl_int nelems)
 {
     if (!is_extension_available(device, "cl_khr_semaphore"))
     {
@@ -210,7 +212,7 @@ int MakeAndRunTest(cl_device_id device, cl_context context,
     cl_int status = TEST_PASS;
     try
     {
-        auto test_fixture = T(device, context, queue);
+        auto test_fixture = T(device, context, queue, nelems);
         status = test_fixture.Run();
     } catch (const std::runtime_error &e)
     {

test_conformance/extensions/cl_khr_semaphore/test_semaphores.cpp

@@ -1,5 +1,5 @@
 //
-// Copyright (c) 2023 The Khronos Group Inc.
+// Copyright (c) 2024 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.
@@ -29,8 +29,8 @@ const char* source = "__kernel void empty() {}";
 struct SimpleSemaphore1 : public SemaphoreTestBase
 {
     SimpleSemaphore1(cl_device_id device, cl_context context,
-                     cl_command_queue queue)
+                     cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -79,8 +79,8 @@ struct SimpleSemaphore1 : public SemaphoreTestBase
 struct SimpleSemaphore2 : public SemaphoreTestBase
 {
     SimpleSemaphore2(cl_device_id device, cl_context context,
-                     cl_command_queue queue)
+                     cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -160,8 +160,8 @@ struct SimpleSemaphore2 : public SemaphoreTestBase
 struct SemaphoreReuse : public SemaphoreTestBase
 {
     SemaphoreReuse(cl_device_id device, cl_context context,
-                   cl_command_queue queue)
+                   cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -250,101 +250,12 @@ struct SemaphoreReuse : public SemaphoreTestBase
     }
 };
 
-template <bool in_order> struct SemaphoreCrossQueue : public SemaphoreTestBase
-{
-    SemaphoreCrossQueue(cl_device_id device, cl_context context,
-                        cl_command_queue queue)
-        : SemaphoreTestBase(device, context, queue)
-    {}
-
-    // Helper function that signals and waits on semaphore across two different
-    // queues.
-    int semaphore_cross_queue_helper(cl_device_id deviceID, cl_context context,
-                                     cl_command_queue queue_1,
-                                     cl_command_queue queue_2)
-    {
-        cl_int err = CL_SUCCESS;
-        // Create semaphore
-        cl_semaphore_properties_khr sema_props[] = {
-            static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_KHR),
-            static_cast<cl_semaphore_properties_khr>(
-                CL_SEMAPHORE_TYPE_BINARY_KHR),
-            0
-        };
-        semaphore =
-            clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err);
-        test_error(err, "Could not create semaphore");
-
-        // Signal semaphore on queue_1
-        clEventWrapper signal_event;
-        err = clEnqueueSignalSemaphoresKHR(queue_1, 1, semaphore, nullptr, 0,
-                                           nullptr, &signal_event);
-        test_error(err, "Could not signal semaphore");
-
-        // Wait semaphore on queue_2
-        clEventWrapper wait_event;
-        err = clEnqueueWaitSemaphoresKHR(queue_2, 1, semaphore, nullptr, 0,
-                                         nullptr, &wait_event);
-        test_error(err, "Could not wait semaphore");
-
-        // Finish queue_1 and queue_2
-        err = clFinish(queue_1);
-        test_error(err, "Could not finish queue");
-
-        err = clFinish(queue_2);
-        test_error(err, "Could not finish queue");
-
-        // Ensure all events are completed
-        test_assert_event_complete(signal_event);
-        test_assert_event_complete(wait_event);
-
-        return TEST_PASS;
-    }
-
-    cl_int run_in_order()
-    {
-        cl_int err = CL_SUCCESS;
-        // Create in-order queues
-        clCommandQueueWrapper queue_1 =
-            clCreateCommandQueue(context, device, 0, &err);
-        test_error(err, "Could not create command queue");
-
-        clCommandQueueWrapper queue_2 =
-            clCreateCommandQueue(context, device, 0, &err);
-        test_error(err, "Could not create command queue");
-
-        return semaphore_cross_queue_helper(device, context, queue_1, queue_2);
-    }
-
-    cl_int run_out_of_order()
-    {
-        cl_int err = CL_SUCCESS;
-        // Create ooo queues
-        clCommandQueueWrapper queue_1 = clCreateCommandQueue(
-            context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
-        test_error(err, "Could not create command queue");
-
-        clCommandQueueWrapper queue_2 = clCreateCommandQueue(
-            context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
-        test_error(err, "Could not create command queue");
-
-        return semaphore_cross_queue_helper(device, context, queue_1, queue_2);
-    }
-
-    cl_int Run() override
-    {
-        if (in_order)
-            return run_in_order();
-        else
-            return run_out_of_order();
-    }
-};
-
 struct SemaphoreMultiSignal : public SemaphoreTestBase
 {
     SemaphoreMultiSignal(cl_device_id device, cl_context context,
-                         cl_command_queue queue)
+                         cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue), semaphore_second(this)
+        : SemaphoreTestBase(device, context, queue, nelems),
+          semaphore_second(this)
     {}
 
     cl_int Run() override
@@ -406,8 +317,9 @@ struct SemaphoreMultiSignal : public SemaphoreTestBase
 struct SemaphoreMultiWait : public SemaphoreTestBase
 {
     SemaphoreMultiWait(cl_device_id device, cl_context context,
-                       cl_command_queue queue)
+                       cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue), semaphore_second(this)
+        : SemaphoreTestBase(device, context, queue, nelems),
+          semaphore_second(this)
     {}
 
     cl_int Run() override
@@ -469,8 +381,9 @@ struct SemaphoreMultiWait : public SemaphoreTestBase
 struct SemaphoreImportExportFD : public SemaphoreTestBase
 {
     SemaphoreImportExportFD(cl_device_id device, cl_context context,
-                            cl_command_queue queue)
+                            cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue), semaphore_second(this)
+        : SemaphoreTestBase(device, context, queue, nelems),
+          semaphore_second(this)
     {}
 
     cl_int Run() override
@@ -560,7 +473,8 @@ struct SemaphoreImportExportFD : public SemaphoreTestBase
 int test_semaphores_simple_1(cl_device_id deviceID, cl_context context,
                              cl_command_queue defaultQueue, int num_elements)
 {
-    return MakeAndRunTest<SimpleSemaphore1>(deviceID, context, defaultQueue);
+    return MakeAndRunTest<SimpleSemaphore1>(deviceID, context, defaultQueue,
+                                            num_elements);
 }
 
 // Confirm that signal a semaphore with no event dependencies will not result
@@ -568,32 +482,16 @@ int test_semaphores_simple_1(cl_device_id deviceID, cl_context context,
 int test_semaphores_simple_2(cl_device_id deviceID, cl_context context,
                              cl_command_queue defaultQueue, int num_elements)
 {
-    return MakeAndRunTest<SimpleSemaphore2>(deviceID, context, defaultQueue);
+    return MakeAndRunTest<SimpleSemaphore2>(deviceID, context, defaultQueue,
+                                            num_elements);
 }
 
 // Confirm that a semaphore can be reused multiple times
 int test_semaphores_reuse(cl_device_id deviceID, cl_context context,
                           cl_command_queue defaultQueue, int num_elements)
 {
-    return MakeAndRunTest<SemaphoreReuse>(deviceID, context, defaultQueue);
+    return MakeAndRunTest<SemaphoreReuse>(deviceID, context, defaultQueue,
-}
+                                          num_elements);
-
-// Confirm that a semaphore works across different ooo queues
-int test_semaphores_cross_queues_ooo(cl_device_id deviceID, cl_context context,
-                                     cl_command_queue defaultQueue,
-                                     int num_elements)
-{
-    return MakeAndRunTest<SemaphoreCrossQueue<false>>(deviceID, context,
-                                                      defaultQueue);
-}
-
-// Confirm that a semaphore works across different in-order queues
-int test_semaphores_cross_queues_io(cl_device_id deviceID, cl_context context,
-                                    cl_command_queue defaultQueue,
-                                    int num_elements)
-{
-    return MakeAndRunTest<SemaphoreCrossQueue<true>>(deviceID, context,
-                                                     defaultQueue);
 }
 
 // Confirm that we can signal multiple semaphores with one command
@@ -601,15 +499,16 @@ int test_semaphores_multi_signal(cl_device_id deviceID, cl_context context,
                                  cl_command_queue defaultQueue,
                                  int num_elements)
 {
-    return MakeAndRunTest<SemaphoreMultiSignal>(deviceID, context,
+    return MakeAndRunTest<SemaphoreMultiSignal>(deviceID, context, defaultQueue,
-                                                defaultQueue);
+                                                num_elements);
 }
 
 // Confirm that we can wait for multiple semaphores with one command
 int test_semaphores_multi_wait(cl_device_id deviceID, cl_context context,
                                cl_command_queue defaultQueue, int num_elements)
 {
-    return MakeAndRunTest<SemaphoreMultiWait>(deviceID, context, defaultQueue);
+    return MakeAndRunTest<SemaphoreMultiWait>(deviceID, context, defaultQueue,
+                                              num_elements);
 }
 
 // Test it is possible to export a semaphore to a sync fd and import the same
@@ -619,5 +518,5 @@ int test_semaphores_import_export_fd(cl_device_id deviceID, cl_context context,
                                      int num_elements)
 {
     return MakeAndRunTest<SemaphoreImportExportFD>(deviceID, context,
-                                                   defaultQueue);
+                                                   defaultQueue, num_elements);
 }

test_conformance/extensions/cl_khr_semaphore/test_semaphores_cross_queue.cpp

@@ -0,0 +1,343 @@
+//
+// Copyright (c) 2024 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 <vector>
+
+#include "semaphore_base.h"
+
+namespace {
+
+template <bool in_order> struct SemaphoreCrossQueue : public SemaphoreTestBase
+{
+    SemaphoreCrossQueue(cl_device_id device, cl_context context,
+                        cl_command_queue queue, cl_int nelems)
+        : SemaphoreTestBase(device, context, queue, nelems)
+    {}
+
+    // Helper function that signals and waits on semaphore across two different
+    // queues.
+    int semaphore_cross_queue_helper(cl_device_id deviceID, cl_context context,
+                                     cl_command_queue queue_1,
+                                     cl_command_queue queue_2)
+    {
+        cl_int err = CL_SUCCESS;
+        // Create semaphore
+        cl_semaphore_properties_khr sema_props[] = {
+            static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_KHR),
+            static_cast<cl_semaphore_properties_khr>(
+                CL_SEMAPHORE_TYPE_BINARY_KHR),
+            0
+        };
+        semaphore =
+            clCreateSemaphoreWithPropertiesKHR(context, sema_props, &err);
+        test_error(err, "Could not create semaphore");
+
+        // Signal semaphore on queue_1
+        clEventWrapper signal_event;
+        err = clEnqueueSignalSemaphoresKHR(queue_1, 1, semaphore, nullptr, 0,
+                                           nullptr, &signal_event);
+        test_error(err, "Could not signal semaphore");
+
+        // Wait semaphore on queue_2
+        clEventWrapper wait_event;
+        err = clEnqueueWaitSemaphoresKHR(queue_2, 1, semaphore, nullptr, 0,
+                                         nullptr, &wait_event);
+        test_error(err, "Could not wait semaphore");
+
+        // Finish queue_1 and queue_2
+        err = clFinish(queue_1);
+        test_error(err, "Could not finish queue");
+
+        err = clFinish(queue_2);
+        test_error(err, "Could not finish queue");
+
+        // Ensure all events are completed
+        test_assert_event_complete(signal_event);
+        test_assert_event_complete(wait_event);
+
+        return TEST_PASS;
+    }
+
+    cl_int run_in_order()
+    {
+        cl_int err = CL_SUCCESS;
+        // Create in-order queues
+        clCommandQueueWrapper queue_1 =
+            clCreateCommandQueue(context, device, 0, &err);
+        test_error(err, "Could not create command queue");
+
+        clCommandQueueWrapper queue_2 =
+            clCreateCommandQueue(context, device, 0, &err);
+        test_error(err, "Could not create command queue");
+
+        return semaphore_cross_queue_helper(device, context, queue_1, queue_2);
+    }
+
+    cl_int run_out_of_order()
+    {
+        cl_int err = CL_SUCCESS;
+        // Create ooo queues
+        clCommandQueueWrapper queue_1 = clCreateCommandQueue(
+            context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
+        test_error(err, "Could not create command queue");
+
+        clCommandQueueWrapper queue_2 = clCreateCommandQueue(
+            context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
+        test_error(err, "Could not create command queue");
+
+        return semaphore_cross_queue_helper(device, context, queue_1, queue_2);
+    }
+
+    cl_int Run() override
+    {
+        if (in_order)
+            return run_in_order();
+        else
+            return run_out_of_order();
+    }
+};
+
+template <bool single_queue>
+struct SemaphoreOutOfOrderOps : public SemaphoreTestBase
+{
+    SemaphoreOutOfOrderOps(cl_device_id device, cl_context context,
+                           cl_command_queue queue, cl_int nelems)
+        : SemaphoreTestBase(device, context, queue, nelems)
+    {}
+
+    bool SetUp()
+    {
+        cl_int error = CL_SUCCESS;
+
+        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");
+
+        kernel = clCreateKernel(program, "copy", &error);
+        test_error(error, "Failed to create copy kernel");
+
+        // create producer/consumer out-of-order queues
+        producer_queue = clCreateCommandQueue(
+            context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &error);
+        test_error(error, "Could not create command queue");
+
+        if (single_queue)
+        {
+            consumer_queue = producer_queue;
+        }
+        else
+        {
+            consumer_queue = clCreateCommandQueue(
+                context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,
+                &error);
+            test_error(error, "Could not create command queue");
+        }
+
+        // Create semaphore
+        cl_semaphore_properties_khr sema_props[] = {
+            static_cast<cl_semaphore_properties_khr>(CL_SEMAPHORE_TYPE_KHR),
+            static_cast<cl_semaphore_properties_khr>(
+                CL_SEMAPHORE_TYPE_BINARY_KHR),
+            0
+        };
+        semaphore =
+            clCreateSemaphoreWithPropertiesKHR(context, sema_props, &error);
+        test_error(error, "Could not create semaphore");
+
+        // create memory resources
+        in_mem_A = clCreateBuffer(context, CL_MEM_READ_ONLY,
+                                  sizeof(cl_int) * num_elems, nullptr, &error);
+        test_error(error, "clCreateBuffer failed");
+
+        in_mem_B = clCreateBuffer(context, CL_MEM_READ_ONLY,
+                                  sizeof(cl_int) * num_elems, nullptr, &error);
+        test_error(error, "clCreateBuffer failed");
+
+
+        out_mem_A = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
+                                   sizeof(cl_int) * num_elems, nullptr, &error);
+        test_error(error, "clCreateBuffer failed");
+
+        out_mem_B = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
+                                   sizeof(cl_int) * num_elems, nullptr, &error);
+        test_error(error, "clCreateBuffer failed");
+
+        error = clSetKernelArg(kernel, 0, sizeof(in_mem_A), &in_mem_A);
+        test_error(error, "clSetKernelArg failed");
+
+        error = clSetKernelArg(kernel, 1, sizeof(out_mem_A), &out_mem_A);
+        test_error(error, "clSetKernelArg failed");
+
+        return CL_SUCCESS;
+    }
+
+    cl_int Run() override
+    {
+        cl_int err = SetUp();
+        test_error(err, "SetUp failed");
+
+        const cl_int pattern_A = 42;
+        const cl_int pattern_B = 0xACDC;
+
+        // enqueue producer operations
+        err = clEnqueueFillBuffer(producer_queue, in_mem_A, &pattern_A,
+                                  sizeof(cl_int), 0, sizeof(cl_int) * num_elems,
+                                  0, nullptr, nullptr);
+        test_error(err, "clEnqueueReadBuffer failed");
+
+        err = clEnqueueFillBuffer(producer_queue, in_mem_B, &pattern_B,
+                                  sizeof(cl_int), 0, sizeof(cl_int) * num_elems,
+                                  0, nullptr, nullptr);
+        test_error(err, "clEnqueueReadBuffer failed");
+
+        // The semaphore cannot be signaled until the barrier is complete
+        err = clEnqueueBarrierWithWaitList(producer_queue, 0, nullptr, nullptr);
+        test_error(err, " clEnqueueBarrierWithWaitList ");
+
+        if (single_queue)
+        {
+            clEventWrapper sema_wait_event;
+
+            // signal/wait with event dependency
+            err = clEnqueueSignalSemaphoresKHR(producer_queue, 1, semaphore,
+                                               nullptr, 0, nullptr,
+                                               &sema_wait_event);
+            test_error(err, "Could not signal semaphore");
+
+            // consumer and producer queues in sync through wait event
+            err = clEnqueueWaitSemaphoresKHR(consumer_queue, 1, semaphore,
+                                             nullptr, 1, &sema_wait_event,
+                                             nullptr);
+            test_error(err, "Could not wait semaphore");
+        }
+        else
+        {
+            err = clEnqueueSignalSemaphoresKHR(producer_queue, 1, semaphore,
+                                               nullptr, 0, nullptr, nullptr);
+            test_error(err, "Could not signal semaphore");
+
+            err = clEnqueueWaitSemaphoresKHR(consumer_queue, 1, semaphore,
+                                             nullptr, 0, nullptr, nullptr);
+            test_error(err, "Could not wait semaphore");
+        }
+
+        err = clEnqueueBarrierWithWaitList(consumer_queue, 0, nullptr, nullptr);
+        test_error(err, " clEnqueueBarrierWithWaitList ");
+
+        // enqueue consumer operations
+        size_t threads = (size_t)num_elems;
+        err = clEnqueueNDRangeKernel(consumer_queue, kernel, 1, nullptr,
+                                     &threads, nullptr, 0, nullptr, nullptr);
+        test_error(err, "clEnqueueNDRangeKernel failed");
+
+        err = clSetKernelArg(kernel, 0, sizeof(in_mem_B), &in_mem_B);
+        test_error(err, "clSetKernelArg failed");
+
+        err = clSetKernelArg(kernel, 1, sizeof(out_mem_B), &out_mem_B);
+        test_error(err, "clSetKernelArg failed");
+
+        err = clEnqueueNDRangeKernel(consumer_queue, kernel, 1, nullptr,
+                                     &threads, nullptr, 0, nullptr, nullptr);
+        test_error(err, "clEnqueueNDRangeKernel failed");
+
+        err = clEnqueueBarrierWithWaitList(consumer_queue, 0, nullptr, nullptr);
+        test_error(err, " clEnqueueBarrierWithWaitList ");
+
+        std::vector<cl_int> host_buffer(num_elems, 0);
+        auto verify_result = [&](const cl_mem &out_mem, const cl_int pattern) {
+            err = clEnqueueReadBuffer(consumer_queue, out_mem, CL_TRUE, 0,
+                                      sizeof(cl_int) * num_elems,
+                                      host_buffer.data(), 0, nullptr, nullptr);
+            test_error_ret(err, "clEnqueueReadBuffer failed", false);
+
+            for (int i = 0; i < num_elems; i++)
+            {
+                if (pattern != host_buffer[i])
+                {
+                    log_error("Expected %d was %d at index %zu\n", pattern,
+                              host_buffer[i], i);
+                    return false;
+                }
+            }
+            return true;
+        };
+
+        if (!verify_result(out_mem_A, pattern_A)) return TEST_FAIL;
+
+        if (!verify_result(out_mem_B, pattern_B)) return TEST_FAIL;
+
+        return CL_SUCCESS;
+    }
+
+    clKernelWrapper kernel = nullptr;
+    clProgramWrapper program = nullptr;
+    clMemWrapper in_mem_A = nullptr, in_mem_B = nullptr, out_mem_A = nullptr,
+                 out_mem_B = nullptr;
+    clCommandQueueWrapper producer_queue = nullptr;
+    clCommandQueueWrapper consumer_queue = nullptr;
+};
+
+} // anonymous namespace
+
+// Confirm that a semaphore works across different ooo queues
+int test_semaphores_cross_queues_ooo(cl_device_id deviceID, cl_context context,
+                                     cl_command_queue defaultQueue,
+                                     int num_elements)
+{
+    return MakeAndRunTest<SemaphoreCrossQueue<false>>(
+        deviceID, context, defaultQueue, num_elements);
+}
+
+// Confirm that a semaphore works across different in-order queues
+int test_semaphores_cross_queues_io(cl_device_id deviceID, cl_context context,
+                                    cl_command_queue defaultQueue,
+                                    int num_elements)
+{
+    return MakeAndRunTest<SemaphoreCrossQueue<true>>(
+        deviceID, context, defaultQueue, num_elements);
+}
+
+// Confirm that we can synchronize signal/wait commands in single out-of-order
+// queue
+int test_semaphores_ooo_ops_single_queue(cl_device_id deviceID,
+                                         cl_context context,
+                                         cl_command_queue defaultQueue,
+                                         int num_elements)
+{
+    return MakeAndRunTest<SemaphoreOutOfOrderOps<true>>(
+        deviceID, context, defaultQueue, num_elements);
+}
+
+// Confirm that we can synchronize signal/wait commands across two out-of-order
+// queues
+int test_semaphores_ooo_ops_cross_queue(cl_device_id deviceID,
+                                        cl_context context,
+                                        cl_command_queue defaultQueue,
+                                        int num_elements)
+{
+    return MakeAndRunTest<SemaphoreOutOfOrderOps<false>>(
+        deviceID, context, defaultQueue, num_elements);
+}

test_conformance/extensions/cl_khr_semaphore/test_semaphores_negative_create.cpp

@@ -27,8 +27,8 @@ namespace {
 struct CreateInvalidContext : public SemaphoreTestBase
 {
     CreateInvalidContext(cl_device_id device, cl_context context,
-                         cl_command_queue queue)
+                         cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -56,8 +56,8 @@ struct CreateInvalidContext : public SemaphoreTestBase
 struct CreateInvalidProperty : public SemaphoreTestBase
 {
     CreateInvalidProperty(cl_device_id device, cl_context context,
-                          cl_command_queue queue)
+                          cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -130,8 +130,8 @@ struct CreateInvalidProperty : public SemaphoreTestBase
 struct CreateInvalidMultiDeviceProperty : public SemaphoreTestBase
 {
     CreateInvalidMultiDeviceProperty(cl_device_id device, cl_context context,
-                                     cl_command_queue queue)
+                                     cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -204,8 +204,8 @@ struct CreateInvalidMultiDeviceProperty : public SemaphoreTestBase
 struct CreateInvalidDevice : public SemaphoreTestBase
 {
     CreateInvalidDevice(cl_device_id device, cl_context context,
-                        cl_command_queue queue)
+                        cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -312,8 +312,9 @@ struct CreateImportExternalWithInvalidDevice : public SemaphoreTestBase
 {
     CreateImportExternalWithInvalidDevice(cl_device_id device,
                                           cl_context context,
-                                          cl_command_queue queue)
+                                          cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue), semaphore_second(this)
+        : SemaphoreTestBase(device, context, queue, nelems),
+          semaphore_second(this)
     {}
 
     cl_int Run() override
@@ -455,8 +456,8 @@ struct CreateImportExternalWithInvalidDevice : public SemaphoreTestBase
 struct CreateInvalidValue : public SemaphoreTestBase
 {
     CreateInvalidValue(cl_device_id device, cl_context context,
-                       cl_command_queue queue)
+                       cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -517,8 +518,9 @@ struct CreateInvalidValue : public SemaphoreTestBase
 struct CreateInvalidOperation : public SemaphoreTestBase
 {
     CreateInvalidOperation(cl_device_id device, cl_context context,
-                           cl_command_queue queue)
+                           cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue), semaphore_second(this)
+        : SemaphoreTestBase(device, context, queue, nelems),
+          semaphore_second(this)
     {}
 
     cl_int Run() override
@@ -603,7 +605,8 @@ int test_semaphores_negative_create_invalid_context(cl_device_id device,
                                                     cl_command_queue queue,
                                                     int num_elements)
 {
-    return MakeAndRunTest<CreateInvalidContext>(device, context, queue);
+    return MakeAndRunTest<CreateInvalidContext>(device, context, queue,
+                                                num_elements);
 }
 
 // Confirm that creation semaphore with invalid properties return
@@ -613,7 +616,8 @@ int test_semaphores_negative_create_invalid_property(cl_device_id device,
                                                      cl_command_queue queue,
                                                      int num_elements)
 {
-    return MakeAndRunTest<CreateInvalidProperty>(device, context, queue);
+    return MakeAndRunTest<CreateInvalidProperty>(device, context, queue,
+                                                 num_elements);
 }
 
 // Confirm that creation semaphore with multi device property return
@@ -622,8 +626,8 @@ int test_semaphores_negative_create_multi_device_property(
     cl_device_id device, cl_context context, cl_command_queue queue,
     int num_elements)
 {
-    return MakeAndRunTest<CreateInvalidMultiDeviceProperty>(device, context,
+    return MakeAndRunTest<CreateInvalidMultiDeviceProperty>(
-                                                            queue);
+        device, context, queue, num_elements);
 }
 
 // Confirm that creation semaphore with invalid device(s) return
@@ -633,7 +637,8 @@ int test_semaphores_negative_create_invalid_device(cl_device_id device,
                                                    cl_command_queue queue,
                                                    int num_elements)
 {
-    return MakeAndRunTest<CreateInvalidDevice>(device, context, queue);
+    return MakeAndRunTest<CreateInvalidDevice>(device, context, queue,
+                                               num_elements);
 }
 
 // Confirm that creation semaphore with invalid device(s) return
@@ -643,7 +648,7 @@ int test_semaphores_negative_create_import_invalid_device(
     int num_elements)
 {
     return MakeAndRunTest<CreateImportExternalWithInvalidDevice>(
-        device, context, queue);
+        device, context, queue, num_elements);
 }
 
 // Confirm that creation semaphore with invalid props values return
@@ -653,7 +658,8 @@ int test_semaphores_negative_create_invalid_value(cl_device_id device,
                                                   cl_command_queue queue,
                                                   int num_elements)
 {
-    return MakeAndRunTest<CreateInvalidValue>(device, context, queue);
+    return MakeAndRunTest<CreateInvalidValue>(device, context, queue,
+                                              num_elements);
 }
 
 // Confirm that creation semaphore with invalid props values return
@@ -663,5 +669,6 @@ int test_semaphores_negative_create_invalid_operation(cl_device_id device,
                                                       cl_command_queue queue,
                                                       int num_elements)
 {
-    return MakeAndRunTest<CreateInvalidOperation>(device, context, queue);
+    return MakeAndRunTest<CreateInvalidOperation>(device, context, queue,
+                                                  num_elements);
 }

test_conformance/extensions/cl_khr_semaphore/test_semaphores_negative_getinfo.cpp

@@ -25,8 +25,8 @@ namespace {
 struct GetInfoInvalidSemaphore : public SemaphoreTestBase
 {
     GetInfoInvalidSemaphore(cl_device_id device, cl_context context,
-                            cl_command_queue queue)
+                            cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -53,8 +53,8 @@ struct GetInfoInvalidSemaphore : public SemaphoreTestBase
 struct GetInfoInvalidValue : public SemaphoreTestBase
 {
     GetInfoInvalidValue(cl_device_id device, cl_context context,
-                        cl_command_queue queue)
+                        cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -118,7 +118,8 @@ int test_semaphores_negative_get_info_invalid_semaphore(cl_device_id device,
                                                         cl_command_queue queue,
                                                         int num_elements)
 {
-    return MakeAndRunTest<GetInfoInvalidSemaphore>(device, context, queue);
+    return MakeAndRunTest<GetInfoInvalidSemaphore>(device, context, queue,
+                                                   num_elements);
 }
 
 int test_semaphores_negative_get_info_invalid_value(cl_device_id device,
@@ -126,5 +127,6 @@ int test_semaphores_negative_get_info_invalid_value(cl_device_id device,
                                                     cl_command_queue queue,
                                                     int num_elements)
 {
-    return MakeAndRunTest<GetInfoInvalidValue>(device, context, queue);
+    return MakeAndRunTest<GetInfoInvalidValue>(device, context, queue,
+                                               num_elements);
 }

test_conformance/extensions/cl_khr_semaphore/test_semaphores_negative_release_retain.cpp

@@ -29,8 +29,8 @@ namespace {
 struct ReleaseInvalidSemaphore : public SemaphoreTestBase
 {
     ReleaseInvalidSemaphore(cl_device_id device, cl_context context,
-                            cl_command_queue queue)
+                            cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -53,8 +53,8 @@ struct ReleaseInvalidSemaphore : public SemaphoreTestBase
 struct RetainInvalidSemaphore : public SemaphoreTestBase
 {
     RetainInvalidSemaphore(cl_device_id device, cl_context context,
-                           cl_command_queue queue)
+                           cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -79,11 +79,13 @@ struct RetainInvalidSemaphore : public SemaphoreTestBase
 int test_semaphores_negative_release(cl_device_id device, cl_context context,
                                      cl_command_queue queue, int num_elements)
 {
-    return MakeAndRunTest<ReleaseInvalidSemaphore>(device, context, queue);
+    return MakeAndRunTest<ReleaseInvalidSemaphore>(device, context, queue,
+                                                   num_elements);
 }
 
 int test_semaphores_negative_retain(cl_device_id device, cl_context context,
                                     cl_command_queue queue, int num_elements)
 {
-    return MakeAndRunTest<RetainInvalidSemaphore>(device, context, queue);
+    return MakeAndRunTest<RetainInvalidSemaphore>(device, context, queue,
+                                                  num_elements);
 }

test_conformance/extensions/cl_khr_semaphore/test_semaphores_negative_wait_signal.cpp

@@ -57,8 +57,8 @@ struct SubDevicesScopeGuarded
 template <RunMode mode> struct InvalidCommandQueue : public SemaphoreTestBase
 {
     InvalidCommandQueue(cl_device_id device, cl_context context,
-                        cl_command_queue queue)
+                        cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -179,8 +179,8 @@ template <RunMode mode> struct InvalidCommandQueue : public SemaphoreTestBase
 template <RunMode mode> struct InvalidValue : public SemaphoreTestBase
 {
     InvalidValue(cl_device_id device, cl_context context,
-                 cl_command_queue queue)
+                 cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -214,8 +214,8 @@ template <RunMode mode> struct InvalidValue : public SemaphoreTestBase
 template <RunMode mode> struct InvalidSemaphore : public SemaphoreTestBase
 {
     InvalidSemaphore(cl_device_id device, cl_context context,
-                     cl_command_queue queue)
+                     cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -255,8 +255,8 @@ template <RunMode mode> struct InvalidSemaphore : public SemaphoreTestBase
 template <RunMode mode> struct InvalidContext : public SemaphoreTestBase
 {
     InvalidContext(cl_device_id device, cl_context context,
-                   cl_command_queue queue)
+                   cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -343,8 +343,8 @@ template <RunMode mode> struct InvalidContext : public SemaphoreTestBase
 template <RunMode mode> struct InvalidEventWaitList : public SemaphoreTestBase
 {
     InvalidEventWaitList(cl_device_id device, cl_context context,
-                         cl_command_queue queue)
+                         cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -437,8 +437,8 @@ template <RunMode mode> struct InvalidEventWaitList : public SemaphoreTestBase
 template <RunMode mode> struct InvalidEventStatus : public SemaphoreTestBase
 {
     InvalidEventStatus(cl_device_id device, cl_context context,
-                       cl_command_queue queue)
+                       cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -500,7 +500,7 @@ int test_semaphores_negative_wait_invalid_command_queue(cl_device_id device,
                                                         int num_elements)
 {
     return MakeAndRunTest<InvalidCommandQueue<RunMode::RM_WAIT>>(
-        device, context, queue);
+        device, context, queue, num_elements);
 }
 
 int test_semaphores_negative_wait_invalid_value(cl_device_id device,
@@ -509,7 +509,7 @@ int test_semaphores_negative_wait_invalid_value(cl_device_id device,
                                                 int num_elements)
 {
     return MakeAndRunTest<InvalidValue<RunMode::RM_WAIT>>(device, context,
-                                                          queue);
+                                                          queue, num_elements);
 }
 
 int test_semaphores_negative_wait_invalid_semaphore(cl_device_id device,
@@ -517,8 +517,8 @@ int test_semaphores_negative_wait_invalid_semaphore(cl_device_id device,
                                                     cl_command_queue queue,
                                                     int num_elements)
 {
-    return MakeAndRunTest<InvalidSemaphore<RunMode::RM_WAIT>>(device, context,
+    return MakeAndRunTest<InvalidSemaphore<RunMode::RM_WAIT>>(
-                                                              queue);
+        device, context, queue, num_elements);
 }
 
 int test_semaphores_negative_wait_invalid_context(cl_device_id device,
@@ -526,8 +526,8 @@ int test_semaphores_negative_wait_invalid_context(cl_device_id device,
                                                   cl_command_queue queue,
                                                   int num_elements)
 {
-    return MakeAndRunTest<InvalidContext<RunMode::RM_WAIT>>(device, context,
+    return MakeAndRunTest<InvalidContext<RunMode::RM_WAIT>>(
-                                                            queue);
+        device, context, queue, num_elements);
 }
 
 int test_semaphores_negative_wait_invalid_event_wait_list(
@@ -535,7 +535,7 @@ int test_semaphores_negative_wait_invalid_event_wait_list(
     int num_elements)
 {
     return MakeAndRunTest<InvalidEventWaitList<RunMode::RM_WAIT>>(
-        device, context, queue);
+        device, context, queue, num_elements);
 }
 
 int test_semaphores_negative_wait_invalid_event_status(cl_device_id device,
@@ -543,8 +543,8 @@ int test_semaphores_negative_wait_invalid_event_status(cl_device_id device,
                                                        cl_command_queue queue,
                                                        int num_elements)
 {
-    return MakeAndRunTest<InvalidEventStatus<RunMode::RM_WAIT>>(device, context,
+    return MakeAndRunTest<InvalidEventStatus<RunMode::RM_WAIT>>(
-                                                                queue);
+        device, context, queue, num_elements);
 }
 
 int test_semaphores_negative_signal_invalid_command_queue(
@@ -552,7 +552,7 @@ int test_semaphores_negative_signal_invalid_command_queue(
     int num_elements)
 {
     return MakeAndRunTest<InvalidCommandQueue<RunMode::RM_SIGNAL>>(
-        device, context, queue);
+        device, context, queue, num_elements);
 }
 
 int test_semaphores_negative_signal_invalid_value(cl_device_id device,
@@ -560,8 +560,8 @@ int test_semaphores_negative_signal_invalid_value(cl_device_id device,
                                                   cl_command_queue queue,
                                                   int num_elements)
 {
-    return MakeAndRunTest<InvalidValue<RunMode::RM_SIGNAL>>(device, context,
+    return MakeAndRunTest<InvalidValue<RunMode::RM_SIGNAL>>(
-                                                            queue);
+        device, context, queue, num_elements);
 }
 
 int test_semaphores_negative_signal_invalid_semaphore(cl_device_id device,
@@ -569,8 +569,8 @@ int test_semaphores_negative_signal_invalid_semaphore(cl_device_id device,
                                                       cl_command_queue queue,
                                                       int num_elements)
 {
-    return MakeAndRunTest<InvalidSemaphore<RunMode::RM_SIGNAL>>(device, context,
+    return MakeAndRunTest<InvalidSemaphore<RunMode::RM_SIGNAL>>(
-                                                                queue);
+        device, context, queue, num_elements);
 }
 
 int test_semaphores_negative_signal_invalid_context(cl_device_id device,
@@ -578,8 +578,8 @@ int test_semaphores_negative_signal_invalid_context(cl_device_id device,
                                                     cl_command_queue queue,
                                                     int num_elements)
 {
-    return MakeAndRunTest<InvalidContext<RunMode::RM_SIGNAL>>(device, context,
+    return MakeAndRunTest<InvalidContext<RunMode::RM_SIGNAL>>(
-                                                              queue);
+        device, context, queue, num_elements);
 }
 
 int test_semaphores_negative_signal_invalid_event_wait_list(
@@ -587,7 +587,7 @@ int test_semaphores_negative_signal_invalid_event_wait_list(
     int num_elements)
 {
     return MakeAndRunTest<InvalidEventWaitList<RunMode::RM_SIGNAL>>(
-        device, context, queue);
+        device, context, queue, num_elements);
 }
 
 int test_semaphores_negative_signal_invalid_event_status(cl_device_id device,
@@ -596,5 +596,5 @@ int test_semaphores_negative_signal_invalid_event_status(cl_device_id device,
                                                          int num_elements)
 {
     return MakeAndRunTest<InvalidEventStatus<RunMode::RM_SIGNAL>>(
-        device, context, queue);
+        device, context, queue, num_elements);
 }

test_conformance/extensions/cl_khr_semaphore/test_semaphores_queries.cpp

@@ -70,8 +70,8 @@ namespace {
 struct SemaphoreWithDeviceListQueries : public SemaphoreTestBase
 {
     SemaphoreWithDeviceListQueries(cl_device_id device, cl_context context,
-                                   cl_command_queue queue)
+                                   cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -144,8 +144,8 @@ struct SemaphoreWithDeviceListQueries : public SemaphoreTestBase
 struct SemaphoreNoDeviceListQueries : public SemaphoreTestBase
 {
     SemaphoreNoDeviceListQueries(cl_device_id device, cl_context context,
-                                 cl_command_queue queue)
+                                 cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -178,8 +178,8 @@ struct SemaphoreNoDeviceListQueries : public SemaphoreTestBase
 struct SemaphoreMultiDeviceContextQueries : public SemaphoreTestBase
 {
     SemaphoreMultiDeviceContextQueries(cl_device_id device, cl_context context,
-                                       cl_command_queue queue)
+                                       cl_command_queue queue, cl_int nelems)
-        : SemaphoreTestBase(device, context, queue)
+        : SemaphoreTestBase(device, context, queue, nelems)
     {}
 
     cl_int Run() override
@@ -260,8 +260,8 @@ int test_semaphores_device_list_queries(cl_device_id deviceID,
                                         cl_command_queue defaultQueue,
                                         int num_elements)
 {
-    return MakeAndRunTest<SemaphoreWithDeviceListQueries>(deviceID, context,
+    return MakeAndRunTest<SemaphoreWithDeviceListQueries>(
-                                                          defaultQueue);
+        deviceID, context, defaultQueue, num_elements);
 }
 
 // Confirm the semaphore without device list can be successfully queried
@@ -270,8 +270,8 @@ int test_semaphores_no_device_list_queries(cl_device_id deviceID,
                                            cl_command_queue defaultQueue,
                                            int num_elements)
 {
-    return MakeAndRunTest<SemaphoreNoDeviceListQueries>(deviceID, context,
+    return MakeAndRunTest<SemaphoreNoDeviceListQueries>(
-                                                        defaultQueue);
+        deviceID, context, defaultQueue, num_elements);
 }
 
 // Confirm the semaphore created with multi-device context can be successfully
@@ -281,6 +281,6 @@ int test_semaphores_multi_device_context_queries(cl_device_id deviceID,
                                                  cl_command_queue defaultQueue,
                                                  int num_elements)
 {
-    return MakeAndRunTest<SemaphoreMultiDeviceContextQueries>(deviceID, context,
+    return MakeAndRunTest<SemaphoreMultiDeviceContextQueries>(
-                                                              defaultQueue);
+        deviceID, context, defaultQueue, num_elements);
 }