[SYCL] do not store last event as a dendency for in-order queues

igchor · igchor · commit ca630f37f690 · 2025-04-15T23:53:59.000Z
For in-order queues, the ordering is guaranteed by UR. The
only expception is Host Task - if the last event is a host
task, then enqueue a barrier.
diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp
@@ -136,18 +136,19 @@ queue_impl::getExtendDependencyList(const std::vector<event> &DepEvents,
     return DepEvents;
 
   QueueLock.lock();
-  EventImplPtr ExtraEvent = MGraph.expired() ? MDefaultGraphDeps.LastEventPtr
-                                             : MExtGraphDeps.LastEventPtr;
   std::optional<event> ExternalEvent = popExternalEvent();
 
-  if (!ExternalEvent && !ExtraEvent)
+  if (!ExternalEvent && !LastHostTaskEvent)
     return DepEvents;
 
   MutableVec = DepEvents;
   if (ExternalEvent)
     MutableVec.push_back(*ExternalEvent);
-  if (ExtraEvent)
-    MutableVec.push_back(detail::createSyclObjFromImpl<event>(ExtraEvent));
+  if (LastHostTaskEvent) {
+    MutableVec.push_back(
+        detail::createSyclObjFromImpl<event>(LastHostTaskEvent));
+    LastHostTaskEvent = nullptr;
+  }
   return MutableVec;
 }
 
@@ -282,20 +283,21 @@ event queue_impl::memcpyFromDeviceGlobal(
       DeviceGlobalPtr, IsDeviceImageScope, Self, NumBytes, Offset, Dest);
 }
 
-sycl::detail::optional<event> queue_impl::getLastEvent() {
-  // The external event is required to finish last if set, so it is considered
-  // the last event if present.
+sycl::detail::optional<event>
+queue_impl::getLastEvent(const std::shared_ptr<queue_impl> &Self) {
   if (std::optional<event> ExternalEvent = MInOrderExternalEvent.read())
     return ExternalEvent;
 
-  std::lock_guard<std::mutex> Lock{MMutex};
-  if (MGraph.expired() && !MDefaultGraphDeps.LastEventPtr)
+  if (MEmpty) {
     return std::nullopt;
-  if (MDiscardEvents)
-    return createDiscardedEvent();
-  if (!MGraph.expired() && MExtGraphDeps.LastEventPtr)
-    return detail::createSyclObjFromImpl<event>(MExtGraphDeps.LastEventPtr);
-  return detail::createSyclObjFromImpl<event>(MDefaultGraphDeps.LastEventPtr);
+  }
+
+  if (LastHostTaskEvent) {
+    return detail::createSyclObjFromImpl<event>(LastHostTaskEvent);
+  }
+
+  // We insert a marker to represent an event at end.
+  return detail::createSyclObjFromImpl<event>(insertMarkerEvent(Self));
 }
 
 void queue_impl::addEvent(const event &Event) {
@@ -375,11 +377,11 @@ event queue_impl::submitWithHandler(const std::shared_ptr<queue_impl> &Self,
   };
   detail::type_erased_cgfo_ty CGF{L};
 
-  if (!CallerNeedsEvent && supportsDiscardingPiEvents()) {
-    submit_without_event(CGF, Self, SI,
-                         /*CodeLoc*/ {}, /*IsTopCodeLoc*/ true);
-    return createDiscardedEvent();
-  }
+  // if (!CallerNeedsEvent && supportsDiscardingPiEvents()) {
+  //   submit_without_event(CGF, Self, SI,
+  //                        /*CodeLoc*/ {}, /*IsTopCodeLoc*/ true);
+  //   return createDiscardedEvent();
+  // }
   return submit_with_event(CGF, Self, SI,
                            /*CodeLoc*/ {}, /*IsTopCodeLoc*/ true);
 }
@@ -396,6 +398,32 @@ event queue_impl::submitMemOpHelper(const std::shared_ptr<queue_impl> &Self,
   {
     std::unique_lock<std::mutex> Lock(MMutex, std::defer_lock);
 
+    if (isInOrder()) {
+      Lock.lock();
+      std::optional<event> ExternalEvent = popExternalEvent();
+
+      if (LastHostTaskEvent) {
+        // TODO: this should be in finalize?
+        LastHostTaskEvent->wait(LastHostTaskEvent);
+        LastHostTaskEvent = nullptr;
+      }
+
+      std::vector<event> WaitEvents;
+      if (ExternalEvent)
+        WaitEvents.emplace_back(std::move(*ExternalEvent));
+
+      event ResEvent = prepareSYCLEventAssociatedWithQueue(Self);
+      const auto &EventImpl = detail::getSyclObjImpl(ResEvent);
+      {
+        NestedCallsTracker tracker;
+        ur_event_handle_t UREvent = nullptr;
+        MemOpFunc(MemOpArgs..., getUrEvents(WaitEvents), &UREvent, EventImpl);
+        EventImpl->setHandle(UREvent);
+        EventImpl->setEnqueued();
+      }
+      return discard_or_return(ResEvent);
+    }
+
     std::vector<event> MutableDepEvents;
     const std::vector<event> &ExpandedDepEvents =
         getExtendDependencyList(DepEvents, MutableDepEvents, Lock);
@@ -404,22 +432,22 @@ event queue_impl::submitMemOpHelper(const std::shared_ptr<queue_impl> &Self,
     // handler rather than by-passing the scheduler.
     if (MGraph.expired() && Scheduler::areEventsSafeForSchedulerBypass(
                                 ExpandedDepEvents, MContext)) {
-      if ((MDiscardEvents || !CallerNeedsEvent) &&
-          supportsDiscardingPiEvents()) {
-        NestedCallsTracker tracker;
-        MemOpFunc(MemOpArgs..., getUrEvents(ExpandedDepEvents),
-                  /*PiEvent*/ nullptr, /*EventImplPtr*/ nullptr);
-
-        event DiscardedEvent = createDiscardedEvent();
-        if (isInOrder()) {
-          // Store the discarded event for proper in-order dependency tracking.
-          auto &EventToStoreIn = MGraph.expired()
-                                     ? MDefaultGraphDeps.LastEventPtr
-                                     : MExtGraphDeps.LastEventPtr;
-          EventToStoreIn = detail::getSyclObjImpl(DiscardedEvent);
-        }
-        return DiscardedEvent;
-      }
+      // if ((MDiscardEvents || !CallerNeedsEvent) &&
+      //     supportsDiscardingPiEvents()) {
+      //   NestedCallsTracker tracker;
+      //   MemOpFunc(MemOpArgs..., getUrEvents(ExpandedDepEvents),
+      //             /*PiEvent*/ nullptr, /*EventImplPtr*/ nullptr);
+
+      //   event DiscardedEvent = createDiscardedEvent();
+      //   if (isInOrder()) {
+      //     // Store the discarded event for proper in-order dependency
+      //     tracking. auto &EventToStoreIn = MGraph.expired()
+      //                                ? MDefaultGraphDeps.LastEventPtr
+      //                                : MExtGraphDeps.LastEventPtr;
+      //     EventToStoreIn = detail::getSyclObjImpl(DiscardedEvent);
+      //   }
+      //   return DiscardedEvent;
+      // }
 
       event ResEvent = prepareSYCLEventAssociatedWithQueue(Self);
       const auto &EventImpl = detail::getSyclObjImpl(ResEvent);
@@ -443,12 +471,6 @@ event queue_impl::submitMemOpHelper(const std::shared_ptr<queue_impl> &Self,
         }
       }
 
-      if (isInOrder()) {
-        auto &EventToStoreIn = MGraph.expired() ? MDefaultGraphDeps.LastEventPtr
-                                                : MExtGraphDeps.LastEventPtr;
-        EventToStoreIn = EventImpl;
-      }
-
       return discard_or_return(ResEvent);
     }
   }
@@ -569,6 +591,11 @@ void queue_impl::wait(const detail::code_location &CodeLoc) {
   std::vector<std::weak_ptr<event_impl>> WeakEvents;
   {
     std::lock_guard<std::mutex> Lock(MMutex);
+    if (LastHostTaskEvent) {
+      LastHostTaskEvent->wait(LastHostTaskEvent);
+      LastHostTaskEvent = nullptr;
+    }
+
     WeakEvents.swap(MEventsWeak);
 
     MMissedCleanupRequests.unset(
@@ -684,23 +711,6 @@ ur_native_handle_t queue_impl::getNative(int32_t &NativeHandleDesc) const {
 }
 
 bool queue_impl::ext_oneapi_empty() const {
-  // If we have in-order queue where events are not discarded then just check
-  // the status of the last event.
-  if (isInOrder() && !MDiscardEvents) {
-    std::lock_guard<std::mutex> Lock(MMutex);
-    // If there is no last event we know that no work has been submitted, so it
-    // must be trivially empty.
-    if (!MDefaultGraphDeps.LastEventPtr)
-      return true;
-    // Otherwise, check if the last event is finished.
-    // Note that we fall back to the backend query if the event was discarded,
-    // which may happend despite the queue not being a discard event queue.
-    if (!MDefaultGraphDeps.LastEventPtr->isDiscarded())
-      return MDefaultGraphDeps.LastEventPtr
-                 ->get_info<info::event::command_execution_status>() ==
-             info::event_command_status::complete;
-  }
-
   // Check the status of the backend queue if this is not a host queue.
   ur_bool_t IsReady = false;
   getAdapter()->call<UrApiKind::urQueueGetInfo>(
diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp
@@ -180,8 +180,6 @@ class queue_impl {
 #endif
   }
 
-  sycl::detail::optional<event> getLastEvent();
-
 private:
   void queue_impl_interop(ur_queue_handle_t UrQueue) {
     if (has_property<ext::oneapi::property::queue::discard_events>() &&
@@ -676,6 +674,9 @@ class queue_impl {
     return Result;
   }
 
+  sycl::detail::optional<event>
+  getLastEvent(const std::shared_ptr<queue_impl> &Self);
+
   const std::vector<event> &
   getExtendDependencyList(const std::vector<event> &DepEvents,
                           std::vector<event> &MutableVec,
@@ -742,44 +743,32 @@ class queue_impl {
     // Hence, here is the lock for thread-safety.
     std::lock_guard<std::mutex> Lock{MMutex};
 
-    auto &EventToBuildDeps = MGraph.expired() ? MDefaultGraphDeps.LastEventPtr
-                                              : MExtGraphDeps.LastEventPtr;
-
-    // This dependency is needed for the following purposes:
-    //    - host tasks are handled by the runtime and cannot be implicitly
-    //    synchronized by the backend.
-    //    - to prevent the 2nd kernel enqueue when the 1st kernel is blocked
-    //    by a host task. This dependency allows to build the enqueue order in
-    //    the RT but will not be passed to the backend. See getPIEvents in
-    //    Command.
-    if (EventToBuildDeps) {
-      // In the case where the last event was discarded and we are to run a
-      // host_task, we insert a barrier into the queue and use the resulting
-      // event as the dependency for the host_task.
-      // Note that host_task events can never be discarded, so this will not
-      // insert barriers between host_task enqueues.
-      if (EventToBuildDeps->isDiscarded() &&
-          Handler.getType() == CGType::CodeplayHostTask)
-        EventToBuildDeps = insertHelperBarrier(Handler);
-
-      // depends_on after an async alloc is explicitly disallowed. Async alloc
-      // handles in order queue dependencies preemptively, so we skip them.
-      // Note: This could be improved by moving the handling of dependencies
-      // to before calling the CGF.
-      if (!EventToBuildDeps->isDiscarded() &&
-          !(Handler.getType() == CGType::AsyncAlloc))
-        Handler.depends_on(EventToBuildDeps);
-    }
-
     // If there is an external event set, add it as a dependency and clear it.
     // We do not need to hold the lock as MLastEventMtx will ensure the last
     // event reflects the corresponding external event dependence as well.
     std::optional<event> ExternalEvent = popExternalEvent();
     if (ExternalEvent)
       Handler.depends_on(*ExternalEvent);
 
+    if (LastHostTaskEvent && Handler.getType() == CGType::CodeplayHostTask) {
+      // is this needed?
+      Handler.depends_on(
+          detail::createSyclObjFromImpl<event>(LastHostTaskEvent));
+      LastHostTaskEvent = nullptr;
+    } else if (!LastHostTaskEvent &&
+               Handler.getType() == CGType::CodeplayHostTask) {
+      auto Event = insertHelperBarrier(Handler);
+      Handler.depends_on(Event);
+    } else if (LastHostTaskEvent) {
+      LastHostTaskEvent->wait(LastHostTaskEvent);
+      LastHostTaskEvent = nullptr;
+    }
+
     auto EventRet = Handler.finalize();
-    EventToBuildDeps = getSyclObjImpl(EventRet);
+
+    if (getSyclObjImpl(EventRet)->isHost()) {
+      LastHostTaskEvent = getSyclObjImpl(EventRet);
+    }
 
     return EventRet;
   }
@@ -849,6 +838,7 @@ class queue_impl {
   template <typename HandlerType = handler>
   event finalizeHandler(HandlerType &Handler,
                         const optional<SubmitPostProcessF> &PostProcessorFunc) {
+    MEmpty = false;
     if (PostProcessorFunc) {
       return finalizeHandlerPostProcess(Handler, PostProcessorFunc);
     } else {
@@ -956,6 +946,11 @@ class queue_impl {
   /// List of queues created for FPGA device from a single SYCL queue.
   ur_queue_handle_t MQueue;
 
+  // TODO: this is for in-order queue only. Move it to separate struct.
+  EventImplPtr LastHostTaskEvent = nullptr;
+
+  bool MEmpty = true;
+
   // Access should be guarded with MMutex
   struct DependencyTrackingItems {
     // This event is employed for enhanced dependency tracking with in-order
diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp
@@ -293,24 +293,6 @@ void queue::wait_and_throw_proxy(const detail::code_location &CodeLoc) {
   impl->wait_and_throw(CodeLoc);
 }
 
-static event
-getBarrierEventForInorderQueueHelper(const detail::QueueImplPtr QueueImpl) {
-  // This function should not be called when a queue is recording to a graph,
-  // as a graph can record from multiple queues and we cannot guarantee the
-  // last node added by an in-order queue will be the last node added to the
-  // graph.
-  assert(!QueueImpl->hasCommandGraph() &&
-         "Should not be called in on graph recording.");
-
-  sycl::detail::optional<event> LastEvent = QueueImpl->getLastEvent();
-  if (LastEvent)
-    return *LastEvent;
-
-  // If there was no last event, we create an empty one.
-  return detail::createSyclObjFromImpl<event>(
-      std::make_shared<detail::event_impl>(std::nullopt));
-}
-
 /// Prevents any commands submitted afterward to this queue from executing
 /// until all commands previously submitted to this queue have entered the
 /// complete state.
@@ -321,10 +303,7 @@ getBarrierEventForInorderQueueHelper(const detail::QueueImplPtr QueueImpl) {
 event queue::ext_oneapi_submit_barrier(const detail::code_location &CodeLoc) {
   if (is_in_order() && !impl->hasCommandGraph() && !impl->MDiscardEvents &&
       !impl->MIsProfilingEnabled) {
-    event InOrderLastEvent = getBarrierEventForInorderQueueHelper(impl);
-    // If the last event was discarded, fall back to enqueuing a barrier.
-    if (!detail::getSyclObjImpl(InOrderLastEvent)->isDiscarded())
-      return InOrderLastEvent;
+    return detail::createSyclObjFromImpl<event>(impl->insertMarkerEvent(impl));
   }
 
   return submit([=](handler &CGH) { CGH.ext_oneapi_barrier(); }, CodeLoc);
@@ -346,13 +325,10 @@ event queue::ext_oneapi_submit_barrier(const std::vector<event> &WaitList,
         auto EventImpl = detail::getSyclObjImpl(Event);
         return (EventImpl->isDefaultConstructed() || EventImpl->isNOP()) &&
                !EventImpl->hasCommandGraph();
-      });
+      }); // TODO: is this needed?
   if (is_in_order() && !impl->hasCommandGraph() && !impl->MDiscardEvents &&
       !impl->MIsProfilingEnabled && AllEventsEmptyOrNop) {
-    event InOrderLastEvent = getBarrierEventForInorderQueueHelper(impl);
-    // If the last event was discarded, fall back to enqueuing a barrier.
-    if (!detail::getSyclObjImpl(InOrderLastEvent)->isDiscarded())
-      return InOrderLastEvent;
+    return detail::createSyclObjFromImpl<event>(impl->insertMarkerEvent(impl));
   }
 
   return submit([=](handler &CGH) { CGH.ext_oneapi_barrier(WaitList); },
@@ -432,20 +408,7 @@ sycl::detail::optional<event> queue::ext_oneapi_get_last_event_impl() const {
         make_error_code(errc::invalid),
         "ext_oneapi_get_last_event() can only be called on in-order queues.");
 
-  sycl::detail::optional<event> LastEvent = impl->getLastEvent();
-
-  // If there was no last event, the queue is yet to have any work submitted and
-  // we return a std::nullopt.
-  if (!LastEvent)
-    return std::nullopt;
-
-  // If the last event was discarded or a NOP, we insert a marker to represent
-  // an event at end.
-  auto LastEventImpl = detail::getSyclObjImpl(*LastEvent);
-  if (LastEventImpl->isDiscarded() || LastEventImpl->isNOP())
-    LastEvent =
-        detail::createSyclObjFromImpl<event>(impl->insertMarkerEvent(impl));
-  return LastEvent;
+  return impl->getLastEvent(impl);
 }
 
 void queue::ext_oneapi_set_external_event(const event &external_event) {
diff --git a/sycl/test-e2e/InOrderEventsExt/get_last_event.cpp b/sycl/test-e2e/InOrderEventsExt/get_last_event.cpp
@@ -36,10 +36,6 @@ int Check(const sycl::queue &Q, const char *CheckName, const F &CheckFunc) {
               << std::endl;
     return 1;
   }
-  if (*E != *LastEvent) {
-    std::cout << "Failed " << CheckName << std::endl;
-    return 1;
-  }
   return 0;
 }
 
diff --git a/sycl/test-e2e/InorderQueue/in_order_ext_oneapi_submit_barrier.cpp b/sycl/test-e2e/InorderQueue/in_order_ext_oneapi_submit_barrier.cpp

Original file line number	Diff line number	Diff line change
`@@ -36,10 +36,6 @@ int Check(const sycl::queue &Q, const char *CheckName, const F &CheckFunc) {`
`36`	`36`	`<< std::endl;`
`37`	`37`	`return 1;`
`38`	`38`	`}`
`39`		`- if (E != LastEvent) {`
`40`		`- std::cout << "Failed " << CheckName << std::endl;`
`41`		`- return 1;`
`42`		`- }`
`43`	`39`	`return 0;`
`44`	`40`	`}`
`45`	`41`