[SYCL] Fix depends_on usage with barriers

sycl/source/detail/scheduler/commands.cpp

@@ -3511,6 +3511,18 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
     const AdapterPtr &Adapter = MQueue->getAdapter();
     if (MEvent != nullptr)
       MEvent->setHostEnqueueTime();
+    // User can specify explicit dependencies via depends_on call that we should
+    // honor here. It is very important for cross queue dependencies. We wait
+    // them explicitly since barrier w/o wait list waits for all commands
+    // submitted before and we can't add new dependencies to its wait list.
+    // Output event for wait operation is not requested since barrier is
+    // submitted immediately after and should synchronize it internally.
+    if (RawEvents.size()) {
+      auto Result = Adapter->call_nocheck<UrApiKind::urEnqueueEventsWait>(
+          MQueue->getHandleRef(), RawEvents.size(), &RawEvents[0], nullptr);
+      if (Result != UR_RESULT_SUCCESS)
+        return Result;
+    }
     if (auto Result =
             Adapter->call_nocheck<UrApiKind::urEnqueueEventsWaitWithBarrierExt>(
                 MQueue->getHandleRef(), &Properties, 0, nullptr, Event);
@@ -3545,6 +3557,12 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
     const AdapterPtr &Adapter = MQueue->getAdapter();
     if (MEvent != nullptr)
       MEvent->setHostEnqueueTime();
+    // User can specify explicit dependencies via depends_on call that we should
+    // honor here. It is very important for cross queue dependencies. Adding
+    // them to the barrier wait list since barrier w/ wait list waits only for
+    // the events provided in wait list and we can just extend the list.
+    UrEvents.insert(UrEvents.end(), RawEvents.begin(), RawEvents.end());
+
     if (auto Result =
             Adapter->call_nocheck<UrApiKind::urEnqueueEventsWaitWithBarrierExt>(
                 MQueue->getHandleRef(), &Properties, UrEvents.size(),

sycl/unittests/scheduler/BarrierDependencies.cpp

@@ -0,0 +1,151 @@
+//==-------- BarrierDependencies.cpp --- Scheduler unit tests --------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "SchedulerTest.hpp"
+#include "SchedulerTestUtils.hpp"
+
+#include <helpers/TestKernel.hpp>
+#include <helpers/UrMock.hpp>
+
+#include <detail/event_impl.hpp>
+
+#include <gtest/gtest.h>
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+std::vector<ur_event_handle_t> EventsInWaitList;
+bool EventsWaitVisited = false;
+static ur_result_t redefinedEventWait(void *pParams) {
+  EventsWaitVisited = true;
+
+  auto params = *static_cast<ur_enqueue_events_wait_params_t *>(pParams);
+  for (size_t i = 0; i < *params.pnumEventsInWaitList; ++i)
+    EventsInWaitList.push_back((*params.pphEventWaitList)[i]);
+
+  return UR_RESULT_SUCCESS;
+}
+
+std::vector<ur_event_handle_t> BarrierEventsInWaitList;
+bool BarrierEventsWaitVisited = false;
+ur_result_t redefinedEnqueueEventsWaitWithBarrierExt(void *pParams) {
+  BarrierEventsWaitVisited = true;
+
+  auto params =
+      *static_cast<ur_enqueue_events_wait_with_barrier_ext_params_t *>(pParams);
+  for (auto i = 0u; i < *params.pnumEventsInWaitList; i++) {
+    BarrierEventsInWaitList.push_back((*params.pphEventWaitList)[i]);
+  }
+  return UR_RESULT_SUCCESS;
+}
+
+void clearGlobals() {
+  EventsInWaitList.clear();
+  BarrierEventsInWaitList.clear();
+  BarrierEventsWaitVisited = false;
+  EventsWaitVisited = false;
+}
+
+TEST_F(SchedulerTest, BarrierWithDependsOn) {
+  clearGlobals();
+
+  sycl::unittest::UrMock<> Mock;
+  sycl::platform Plt = sycl::platform();
+  mock::getCallbacks().set_after_callback("urEnqueueEventsWait",
+                                          &redefinedEventWait);
+  mock::getCallbacks().set_after_callback(
+      "urEnqueueEventsWaitWithBarrierExt",
+      &redefinedEnqueueEventsWaitWithBarrierExt);
+
+  context Ctx{Plt};
+  queue QueueA{Ctx, default_selector_v, property::queue::in_order()};
+  queue QueueB{Ctx, default_selector_v, property::queue::in_order()};
+
+  auto EventA =
+      QueueA.submit([&](sycl::handler &h) { h.ext_oneapi_barrier(); });
+  std::shared_ptr<detail::event_impl> EventAImpl =
+      detail::getSyclObjImpl(EventA);
+  // it means that command is enqueued
+  ASSERT_NE(EventAImpl->getHandle(), nullptr);
+
+  ASSERT_FALSE(EventsWaitVisited);
+  ASSERT_TRUE(BarrierEventsWaitVisited);
+  ASSERT_EQ(BarrierEventsInWaitList.size(), 0u);
+
+  clearGlobals();
+  auto EventB = QueueB.submit([&](sycl::handler &h) {
+    h.depends_on(EventA);
+    h.ext_oneapi_barrier();
+  });
+  std::shared_ptr<detail::event_impl> EventBImpl =
+      detail::getSyclObjImpl(EventB);
+  // it means that command is enqueued
+  ASSERT_NE(EventBImpl->getHandle(), nullptr);
+
+  ASSERT_TRUE(EventsWaitVisited);
+  ASSERT_EQ(EventsInWaitList.size(), 1u);
+  EXPECT_EQ(EventsInWaitList[0], EventAImpl->getHandle());
+
+  ASSERT_TRUE(BarrierEventsWaitVisited);
+  ASSERT_EQ(BarrierEventsInWaitList.size(), 0u);
+
+  QueueA.wait();
+  QueueB.wait();
+}
+
+TEST_F(SchedulerTest, BarrierWaitListWithDependsOn) {
+  clearGlobals();
+
+  sycl::unittest::UrMock<> Mock;
+  sycl::platform Plt = sycl::platform();
+  mock::getCallbacks().set_after_callback("urEnqueueEventsWait",
+                                          &redefinedEventWait);
+  mock::getCallbacks().set_after_callback(
+      "urEnqueueEventsWaitWithBarrierExt",
+      &redefinedEnqueueEventsWaitWithBarrierExt);
+
+  context Ctx{Plt};
+  queue QueueA{Ctx, default_selector_v, property::queue::in_order()};
+  queue QueueB{Ctx, default_selector_v, property::queue::in_order()};
+
+  auto EventA =
+      QueueA.submit([&](sycl::handler &h) { h.ext_oneapi_barrier(); });
+  auto EventA2 =
+      QueueA.submit([&](sycl::handler &h) { h.ext_oneapi_barrier(); });
+  std::shared_ptr<detail::event_impl> EventAImpl =
+      detail::getSyclObjImpl(EventA);
+  std::shared_ptr<detail::event_impl> EventA2Impl =
+      detail::getSyclObjImpl(EventA2);
+  // it means that command is enqueued
+  ASSERT_NE(EventAImpl->getHandle(), nullptr);
+  ASSERT_NE(EventA2Impl->getHandle(), nullptr);
+
+  ASSERT_FALSE(EventsWaitVisited);
+  ASSERT_TRUE(BarrierEventsWaitVisited);
+  ASSERT_EQ(BarrierEventsInWaitList.size(), 0u);
+
+  clearGlobals();
+  auto EventB = QueueB.submit([&](sycl::handler &h) {
+    h.depends_on(EventA);
+    h.ext_oneapi_barrier({EventA2});
+  });
+  std::shared_ptr<detail::event_impl> EventBImpl =
+      detail::getSyclObjImpl(EventB);
+  // it means that command is enqueued
+  ASSERT_NE(EventBImpl->getHandle(), nullptr);
+
+  ASSERT_FALSE(EventsWaitVisited);
+  ASSERT_TRUE(BarrierEventsWaitVisited);
+  ASSERT_EQ(BarrierEventsInWaitList.size(), 2u);
+  EXPECT_EQ(BarrierEventsInWaitList[0], EventA2Impl->getHandle());
+  EXPECT_EQ(BarrierEventsInWaitList[1], EventAImpl->getHandle());
+
+  QueueA.wait();
+  QueueB.wait();
+}

sycl/unittests/scheduler/CMakeLists.txt

@@ -21,4 +21,5 @@ add_sycl_unittest(SchedulerTests OBJECT
     EnqueueWithDependsOnDeps.cpp
     AccessorDefaultCtor.cpp
     HostTaskAndBarrier.cpp
+    BarrierDependencies.cpp
 )