[SYCL] Fix in-order queue dependencies for no scheduler path (#15412)

Runtime could submit kernel directly to scheduler if no buffers/streams
are used and if event dependencies are already handled by queue (in case
if it is in-order one). Although check if dependencies are submitted to
the same queue was missed. Now we add events submitted to another queue
but on the same context to event list in kernel launching.

---------

Signed-off-by: Tikhomirova, Kseniya <[email protected]>

Author: KseniyaTikhomirova

sycl/source/detail/scheduler/commands.cpp

@@ -229,7 +229,8 @@ static std::string commandToName(Command::CommandType Type) {
 #endif
 
 std::vector<ur_event_handle_t>
-Command::getUrEvents(const std::vector<EventImplPtr> &EventImpls) const {
+Command::getUrEvents(const std::vector<EventImplPtr> &EventImpls,
+                     const QueueImplPtr &CommandQueue, bool IsHostTaskCommand) {
   std::vector<ur_event_handle_t> RetUrEvents;
   for (auto &EventImpl : EventImpls) {
     auto Handle = EventImpl->getHandle();
@@ -240,8 +241,8 @@ Command::getUrEvents(const std::vector<EventImplPtr> &EventImpls) const {
     // At this stage dependency is definitely ur task and need to check if
     // current one is a host task. In this case we should not skip ur event due
     // to different sync mechanisms for different task types on in-order queue.
-    if (MWorkerQueue && EventImpl->getWorkerQueue() == MWorkerQueue &&
-        MWorkerQueue->isInOrder() && !isHostTask())
+    if (CommandQueue && EventImpl->getWorkerQueue() == CommandQueue &&
+        CommandQueue->isInOrder() && !IsHostTaskCommand)
       continue;
 
     RetUrEvents.push_back(Handle);
@@ -250,6 +251,11 @@ Command::getUrEvents(const std::vector<EventImplPtr> &EventImpls) const {
   return RetUrEvents;
 }
 
+std::vector<ur_event_handle_t>
+Command::getUrEvents(const std::vector<EventImplPtr> &EventImpls) const {
+  return getUrEvents(EventImpls, MWorkerQueue, isHostTask());
+}
+
 // This function is implemented (duplicating getUrEvents a lot) as short term
 // solution for the issue that barrier with wait list could not
 // handle empty ur event handles when kernel is enqueued on host task

sycl/source/detail/scheduler/commands.hpp

@@ -240,6 +240,10 @@ class Command {
   /// in order queue
   std::vector<ur_event_handle_t>
   getUrEvents(const std::vector<EventImplPtr> &EventImpls) const;
+
+  static std::vector<ur_event_handle_t>
+  getUrEvents(const std::vector<EventImplPtr> &EventImpls,
+              const QueueImplPtr &CommandQueue, bool IsHostTaskCommand);
   /// Collect UR events from EventImpls and filter out some of them in case of
   /// in order queue. Does blocking enqueue if event is expected to produce ur
   /// event but has empty native handle.

sycl/source/handler.cpp

@@ -258,8 +258,8 @@ event handler::finalize() {
       // the graph is not changed, then this faster path is used to submit
       // kernel bypassing scheduler and avoiding CommandGroup, Command objects
       // creation.
-
-      std::vector<ur_event_handle_t> RawEvents;
+      std::vector<ur_event_handle_t> RawEvents =
+          detail::Command::getUrEvents(impl->CGData.MEvents, MQueue, false);
       detail::EventImplPtr NewEvent;
 
 #ifdef XPTI_ENABLE_INSTRUMENTATION

sycl/unittests/scheduler/InOrderQueueDeps.cpp

@@ -125,4 +125,68 @@ TEST_F(SchedulerTest, InOrderQueueIsolatedDeps) {
     EXPECT_TRUE(BarrierCalled);
   }
 }
+
+std::vector<size_t> KernelEventListSize;
+
+inline ur_result_t customEnqueueKernelLaunch(void *pParams) {
+  auto params = *static_cast<ur_enqueue_kernel_launch_params_t *>(pParams);
+  KernelEventListSize.push_back(*params.pnumEventsInWaitList);
+  return UR_RESULT_SUCCESS;
+}
+
+TEST_F(SchedulerTest, TwoInOrderQueuesOnSameContext) {
+  KernelEventListSize.clear();
+  sycl::unittest::UrMock<> Mock;
+  mock::getCallbacks().set_before_callback("urEnqueueKernelLaunch",
+                                           &customEnqueueKernelLaunch);
+
+  sycl::platform Plt = sycl::platform();
+
+  context Ctx{Plt};
+  queue InOrderQueueFirst{Ctx, default_selector_v, property::queue::in_order()};
+  queue InOrderQueueSecond{Ctx, default_selector_v,
+                           property::queue::in_order()};
+
+  event EvFirst = InOrderQueueFirst.submit(
+      [&](sycl::handler &CGH) { CGH.single_task<TestKernel<>>([] {}); });
+  std::ignore = InOrderQueueSecond.submit([&](sycl::handler &CGH) {
+    CGH.depends_on(EvFirst);
+    CGH.single_task<TestKernel<>>([] {});
+  });
+
+  InOrderQueueFirst.wait();
+  InOrderQueueSecond.wait();
+
+  ASSERT_EQ(KernelEventListSize.size(), 2u);
+  EXPECT_EQ(KernelEventListSize[0] /*EventsCount*/, 0u);
+  EXPECT_EQ(KernelEventListSize[1] /*EventsCount*/, 1u);
+}
+
+TEST_F(SchedulerTest, InOrderQueueNoSchedulerPath) {
+  KernelEventListSize.clear();
+  sycl::unittest::UrMock<> Mock;
+  mock::getCallbacks().set_before_callback("urEnqueueKernelLaunch",
+                                           &customEnqueueKernelLaunch);
+
+  sycl::platform Plt = sycl::platform();
+
+  context Ctx{Plt};
+  queue InOrderQueue{Ctx, default_selector_v, property::queue::in_order()};
+
+  event EvFirst = InOrderQueue.submit(
+      [&](sycl::handler &CGH) { CGH.single_task<TestKernel<>>([] {}); });
+  std::ignore = InOrderQueue.submit([&](sycl::handler &CGH) {
+    CGH.depends_on(EvFirst);
+    CGH.single_task<TestKernel<>>([] {});
+  });
+
+  InOrderQueue.wait();
+
+  ASSERT_EQ(KernelEventListSize.size(), 2u);
+  EXPECT_EQ(KernelEventListSize[0] /*EventsCount*/, 0u);
+  // native device events for device kernel submitted to the same in-order queue
+  // don't need to be explicitly passed as dependencies
+  EXPECT_EQ(KernelEventListSize[1] /*EventsCount*/, 0u);
+}
+
 } // anonymous namespace