[SYCL] Extend scope of scheduler bypass to safe to bypass events (#16735)

1) Aling scheduler bypass conditions for CG and for memory ops. 
2) Connect an event returned from scheduler bypass case with dependent
events.

---------

Co-authored-by: Sergey Semenov <[email protected]>

Author: Alexandr-Konovalov

sycl/source/detail/queue_impl.cpp

@@ -481,14 +481,25 @@ event queue_impl::submitMemOpHelper(const std::shared_ptr<queue_impl> &Self,
       }
 
       event ResEvent = prepareSYCLEventAssociatedWithQueue(Self);
-      auto EventImpl = detail::getSyclObjImpl(ResEvent);
+      const auto &EventImpl = detail::getSyclObjImpl(ResEvent);
       {
         NestedCallsTracker tracker;
         ur_event_handle_t UREvent = nullptr;
         MemOpFunc(MemOpArgs..., getUrEvents(ExpandedDepEvents), &UREvent,
                   EventImpl);
         EventImpl->setHandle(UREvent);
         EventImpl->setEnqueued();
+        // connect returned event with dependent events
+        if (!isInOrder()) {
+          std::vector<EventImplPtr> &ExpandedDepEventImplPtrs =
+              EventImpl->getPreparedDepsEvents();
+          ExpandedDepEventImplPtrs.reserve(ExpandedDepEvents.size());
+          for (const event &DepEvent : ExpandedDepEvents)
+            ExpandedDepEventImplPtrs.push_back(
+                detail::getSyclObjImpl(DepEvent));
+
+          EventImpl->cleanDepEventsThroughOneLevel();
+        }
       }
 
       if (isInOrder()) {

sycl/source/detail/scheduler/scheduler.cpp

@@ -670,7 +670,7 @@ bool CheckEventReadiness(const ContextImplPtr &Context,
 }
 
 bool Scheduler::areEventsSafeForSchedulerBypass(
-    const std::vector<sycl::event> &DepEvents, ContextImplPtr Context) {
+    const std::vector<sycl::event> &DepEvents, const ContextImplPtr &Context) {
 
   return std::all_of(
       DepEvents.begin(), DepEvents.end(), [&Context](const sycl::event &Event) {
@@ -680,7 +680,7 @@ bool Scheduler::areEventsSafeForSchedulerBypass(
 }
 
 bool Scheduler::areEventsSafeForSchedulerBypass(
-    const std::vector<EventImplPtr> &DepEvents, ContextImplPtr Context) {
+    const std::vector<EventImplPtr> &DepEvents, const ContextImplPtr &Context) {
 
   return std::all_of(DepEvents.begin(), DepEvents.end(),
                      [&Context](const EventImplPtr &SyclEventImplPtr) {

sycl/source/detail/scheduler/scheduler.hpp

@@ -479,10 +479,10 @@ class Scheduler {
 
   static bool
   areEventsSafeForSchedulerBypass(const std::vector<sycl::event> &DepEvents,
-                                  ContextImplPtr Context);
+                                  const ContextImplPtr &Context);
   static bool
   areEventsSafeForSchedulerBypass(const std::vector<EventImplPtr> &DepEvents,
-                                  ContextImplPtr Context);
+                                  const ContextImplPtr &Context);
 
 protected:
   using RWLockT = std::shared_timed_mutex;

sycl/source/handler.cpp

@@ -468,10 +468,8 @@ event handler::finalize() {
     if (MQueue && !impl->MGraph && !impl->MSubgraphNode &&
         !MQueue->hasCommandGraph() && !impl->CGData.MRequirements.size() &&
         !MStreamStorage.size() &&
-        (!impl->CGData.MEvents.size() ||
-         (MQueue->isInOrder() &&
-          detail::Scheduler::areEventsSafeForSchedulerBypass(
-              impl->CGData.MEvents, MQueue->getContextImplPtr())))) {
+        detail::Scheduler::areEventsSafeForSchedulerBypass(
+            impl->CGData.MEvents, MQueue->getContextImplPtr())) {
       // if user does not add a new dependency to the dependency graph, i.e.
       // the graph is not changed, then this faster path is used to submit
       // kernel bypassing scheduler and avoiding CommandGroup, Command objects
@@ -544,6 +542,11 @@ event handler::finalize() {
 
         EnqueueKernel();
         NewEvent->setEnqueued();
+        // connect returned event with dependent events
+        if (!MQueue->isInOrder()) {
+          NewEvent->getPreparedDepsEvents() = impl->CGData.MEvents;
+          NewEvent->cleanDepEventsThroughOneLevel();
+        }
 
         MLastEvent = detail::createSyclObjFromImpl<event>(std::move(NewEvent));
       }

sycl/test-e2e/XPTI/basic_event_collection_linux.cpp

@@ -32,12 +32,6 @@
 // CHECK-DAG:    from_source : false
 // CHECK-DAG:    kernel_name : typeinfo name for main::{lambda(sycl::_V1::handler&)#1}::operator()(sycl::_V1::handler&) const::{lambda()#1}
 // CHECK-DAG:    sycl_device : {{.*}}
-// CHECK:      Node create
-// CHECK-DAG:   queue_id : {{.*}}
-// CHECK-DAG:   kernel_name : virtual_node[{{.*}}]
-// CHECK-NEXT: Edge create
-// CHECK-DAG:   queue_id : {{.*}}
-// CHECK-DAG:   event : {{.*}}
 // CHECK: Task begin
 // CHECK-DAG:    queue_id : {{.*}}
 // CHECK-DAG:    sym_line_no : {{.*}}

sycl/unittests/scheduler/EnqueueWithDependsOnDeps.cpp

@@ -323,6 +323,7 @@ TEST_F(DependsOnTests, ShortcutFunctionWithWaitList) {
                                            &redefinedextUSMEnqueueMemcpy);
   sycl::queue Queue = detail::createSyclObjFromImpl<queue>(QueueDevImpl);
 
+  // Mock up an incomplete host task
   auto HostTaskEvent =
       Queue.submit([&](sycl::handler &cgh) { cgh.host_task([=]() {}); });
   std::shared_ptr<detail::event_impl> HostTaskEventImpl =
@@ -332,6 +333,7 @@ TEST_F(DependsOnTests, ShortcutFunctionWithWaitList) {
   ASSERT_NE(Cmd, nullptr);
   Cmd->MIsBlockable = true;
   Cmd->MEnqueueStatus = detail::EnqueueResultT::SyclEnqueueBlocked;
+  HostTaskEventImpl->setStateIncomplete();
 
   auto SingleTaskEvent = Queue.submit([&](sycl::handler &cgh) {
     cgh.depends_on(HostTaskEvent);
@@ -341,6 +343,8 @@ TEST_F(DependsOnTests, ShortcutFunctionWithWaitList) {
       detail::getSyclObjImpl(SingleTaskEvent);
   EXPECT_EQ(SingleTaskEventImpl->getHandle(), nullptr);
 
+  // make HostTaskEvent completed, so SingleTaskEvent can be enqueued
+  HostTaskEventImpl->setComplete();
   Cmd->MEnqueueStatus = detail::EnqueueResultT::SyclEnqueueSuccess;
   EventsInWaitList.clear();
 
@@ -375,6 +379,7 @@ TEST_F(DependsOnTests, BarrierWithWaitList) {
   ASSERT_NE(Cmd, nullptr);
   Cmd->MIsBlockable = true;
   Cmd->MEnqueueStatus = detail::EnqueueResultT::SyclEnqueueBlocked;
+  HostTaskEventImpl->setStateIncomplete();
 
   auto SingleTaskEvent = Queue.submit([&](sycl::handler &cgh) {
     cgh.depends_on(HostTaskEvent);
@@ -384,6 +389,7 @@ TEST_F(DependsOnTests, BarrierWithWaitList) {
       detail::getSyclObjImpl(SingleTaskEvent);
   EXPECT_EQ(SingleTaskEventImpl->getHandle(), nullptr);
 
+  HostTaskEventImpl->setComplete();
   Cmd->MEnqueueStatus = detail::EnqueueResultT::SyclEnqueueSuccess;
   EventsInWaitList.clear();