[SYCL] Fix barrier submission after command which bypasses the schedu…

…ler (#15697)

Currently we mistakenly don't mark commands like memcpy, copy, fill etc.
as enqueued, this happens because they are enqueued differently
bypassing 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

@@ -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()) {

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