From b6bfafe7af36155d004f8852e0a362241de128dd Mon Sep 17 00:00:00 2001
From: Artur Gainullin <artur.gainullin@intel.com>
Date: Mon, 14 Oct 2024 22:10:08 -0700
Subject: [PATCH] [SYCL] Fix barrier submission after command which bypasses
 the scheduler

Currently we mistakenly don't mark commands like memcpy, copy, fill etc. as
enqueued, this happens because they are enqueued differently and bypass
the scheduler.

Problem is that if barrier is submitted depending on such event then it
is just omitted.

So, fix the problem by marking such commands as enqueued.
---
 sycl/source/detail/queue_impl.cpp             |  1 +
 sycl/unittests/scheduler/InOrderQueueDeps.cpp | 31 +++++++++++++++++++
 2 files changed, 32 insertions(+)

diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp
index b47cc58689b1a..428f06ea0aaa4 100644
--- a/sycl/source/detail/queue_impl.cpp
+++ b/sycl/source/detail/queue_impl.cpp
@@ -455,6 +455,7 @@ event queue_impl::submitMemOpHelper(const std::shared_ptr<queue_impl> &Self,
         MemOpFunc(MemOpArgs..., getUrEvents(ExpandedDepEvents), &UREvent,
                   EventImpl);
         EventImpl->setHandle(UREvent);
+        EventImpl->setEnqueued();
       }
 
       if (isInOrder()) {
diff --git a/sycl/unittests/scheduler/InOrderQueueDeps.cpp b/sycl/unittests/scheduler/InOrderQueueDeps.cpp
index 0da904bb10602..8a26849fc6255 100644
--- a/sycl/unittests/scheduler/InOrderQueueDeps.cpp
+++ b/sycl/unittests/scheduler/InOrderQueueDeps.cpp
@@ -11,6 +11,7 @@
 
 #include <helpers/TestKernel.hpp>
 #include <helpers/UrMock.hpp>
+#include <sycl/usm.hpp>
 
 #include <iostream>
 #include <memory>
@@ -189,4 +190,34 @@ TEST_F(SchedulerTest, InOrderQueueNoSchedulerPath) {
   EXPECT_EQ(KernelEventListSize[1] /*EventsCount*/, 0u);
 }
 
+// Test that barrier is not filtered out when waitlist contains an event
+// produced by command which is bypassing the scheduler.
+TEST_F(SchedulerTest, BypassSchedulerWithBarrier) {
+  sycl::unittest::UrMock<> Mock;
+  sycl::platform Plt = sycl::platform();
+
+  mock::getCallbacks().set_before_callback(
+      "urEnqueueEventsWaitWithBarrier", &redefinedEnqueueEventsWaitWithBarrier);
+  BarrierCalled = false;
+
+  context Ctx{Plt};
+  queue Q1{Ctx, default_selector_v, property::queue::in_order()};
+  queue Q2{Ctx, default_selector_v, property::queue::in_order()};
+  static constexpr size_t Size = 10;
+
+  int *X = malloc_host<int>(Size, Ctx);
+
+  // Submit a command which bypasses the scheduler.
+  auto FillEvent = Q2.memset(X, 0, sizeof(int) * Size);
+  // Submit a barrier which depends on that event.
+  ExpectedEvent = detail::getSyclObjImpl(FillEvent)->getHandle();
+  auto BarrierQ1 = Q1.ext_oneapi_submit_barrier({FillEvent});
+  Q1.wait();
+  Q2.wait();
+  // Verify that barrier is not filtered out.
+  EXPECT_EQ(BarrierCalled, true);
+
+  free(X, Ctx);
+}
+
 } // anonymous namespace