[SYCL] Do not store last event for in-order queues

unless Host Tasks are used.

Without Host Tasks, we can just rely on UR for ordering.
Having no last event means that ext_oneapi_get_last_event()
needs to submit a barrier to return an event to the user.
Similarly, ext_oneapi_submit_barrier() now always submits
a barrier, even for in-order queues.

Whenever Host Tasks are used we need to start recording
all events. This is needed because of how kernel submission
synchronizes with Host Tasks. With a following scenario:

q.host_task();
q.submit_kernel();
q.host_task():

The kernel won't even be submitted to UR until the first
Host Task completes. To properly synchronize the second
Host Task we need to keep the event describing kernel submission.

Author: igchor

sycl/source/detail/queue_impl.cpp

@@ -282,19 +282,26 @@ event queue_impl::memcpyFromDeviceGlobal(
       DeviceGlobalPtr, IsDeviceImageScope, Self, NumBytes, Offset, Dest);
 }
 
-sycl::detail::optional<event> queue_impl::getLastEvent() {
+sycl::detail::optional<event>
+queue_impl::getLastEvent(const std::shared_ptr<queue_impl> &Self) {
   // The external event is required to finish last if set, so it is considered
   // the last event if present.
   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 (MNoEventMode) {
+    assert(MGraph.expired());
+    assert(!MDefaultGraphDeps.LastEventPtr);
+
+    // We insert a marker to represent an event at end.
+    return detail::createSyclObjFromImpl<event>(insertMarkerEvent(Self));
+  }
   if (!MGraph.expired() && MExtGraphDeps.LastEventPtr)
     return detail::createSyclObjFromImpl<event>(MExtGraphDeps.LastEventPtr);
+  assert(MDefaultGraphDeps.LastEventPtr);
   return detail::createSyclObjFromImpl<event>(MDefaultGraphDeps.LastEventPtr);
 }
 
@@ -339,10 +346,27 @@ event queue_impl::submit_impl(const detail::type_erased_cgfo_ty &CGF,
 
   HandlerImpl->MEventMode = SubmitInfo.EventMode();
 
-  auto Event = finalizeHandler(Handler, SubmitInfo.PostProcessorFunc());
+  auto isInNoEventsMode =
+      isInOrder() && MNoEventMode.load(std::memory_order_relaxed);
+  auto needToStoreEvent = Type == CGType::CodeplayHostTask ||
+                          Streams.size() > 0 || SubmitInfo.PostProcessorFunc();
+
+  if (isInNoEventsMode && !needToStoreEvent) {
+    std::unique_lock<std::mutex> Lock(MMutex);
+
+    // Check the condition again, under the lock to ensure that the
+    // there was no concurrent submit that changed the state.
+    if (MNoEventMode.load(std::memory_order_relaxed))
+      return finalizeHandlerInOrderNoEventsUnlocked(Handler);
+  }
 
+  auto Event = finalizeHandler(Handler, SubmitInfo.PostProcessorFunc());
   addEvent(Event);
 
+  // If we taken this path, it's because we need to store the last event.
+  if (isInOrder())
+    assert(!MNoEventMode);
+
   const auto &EventImpl = detail::getSyclObjImpl(Event);
   for (auto &Stream : Streams) {
     // We don't want stream flushing to be blocking operation that is why submit
@@ -460,21 +484,12 @@ 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()) {
+      if (!CallerNeedsEvent && MNoEventMode) {
         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;
+        return createDiscardedEvent();
       }
 
       event ResEvent = prepareSYCLEventAssociatedWithQueue(Self);
@@ -500,7 +515,7 @@ event queue_impl::submitMemOpHelper(const std::shared_ptr<queue_impl> &Self,
         }
       }
 
-      if (isInOrder()) {
+      if (!MNoEventMode) {
         auto &EventToStoreIn = MGraph.expired() ? MDefaultGraphDeps.LastEventPtr
                                                 : MExtGraphDeps.LastEventPtr;
         EventToStoreIn = EventImpl;
@@ -745,14 +760,11 @@ bool queue_impl::ext_oneapi_empty() const {
   // 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.
+    assert((MDefaultGraphDeps.LastEventPtr == nullptr) == MNoEventMode);
     // 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())
+    if (MDefaultGraphDeps.LastEventPtr &&
+        !MDefaultGraphDeps.LastEventPtr->isDiscarded())
       return MDefaultGraphDeps.LastEventPtr
                  ->get_info<info::event::command_execution_status>() ==
              info::event_command_status::complete;
@@ -765,6 +777,11 @@ bool queue_impl::ext_oneapi_empty() const {
   if (!IsReady)
     return false;
 
+  // If got here, it means that LastEventPtr is nullptr (so no possible Host
+  // Tasks) and there is nothing executing on the device.
+  if (isInOrder())
+    return true;
+
   // We may have events like host tasks which are not submitted to the backend
   // queue so we need to get their status separately.
   std::lock_guard<std::mutex> Lock(MMutex);

sycl/source/detail/queue_impl.hpp

@@ -180,7 +180,8 @@ class queue_impl {
 #endif
   }
 
-  sycl::detail::optional<event> getLastEvent();
+  sycl::detail::optional<event>
+  getLastEvent(const std::shared_ptr<queue_impl> &Self);
 
 private:
   void queue_impl_interop(ur_queue_handle_t UrQueue) {
@@ -630,6 +631,7 @@ class queue_impl {
     std::lock_guard<std::mutex> Lock(MMutex);
     MGraph = Graph;
     MExtGraphDeps.reset();
+    MNoEventMode = false;
   }
 
   std::shared_ptr<ext::oneapi::experimental::detail::graph_impl>
@@ -719,6 +721,29 @@ class queue_impl {
     return ResEvent;
   }
 
+  template <typename HandlerType = handler>
+  void synchronizeWithExternalEvent(HandlerType &Handler) {
+    // 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);
+  }
+
+  template <typename HandlerType = handler>
+  event finalizeHandlerInOrderNoEventsUnlocked(HandlerType &Handler) {
+    assert(isInOrder());
+    assert(MGraph.expired());
+    assert(MNoEventMode);
+
+    MEmpty = false;
+
+    synchronizeWithExternalEvent(Handler);
+
+    return Handler.finalize();
+  }
+
   template <typename HandlerType = handler>
   event finalizeHandlerInOrder(HandlerType &Handler) {
     // Accessing and changing of an event isn't atomic operation.
@@ -728,6 +753,16 @@ class queue_impl {
     auto &EventToBuildDeps = MGraph.expired() ? MDefaultGraphDeps.LastEventPtr
                                               : MExtGraphDeps.LastEventPtr;
 
+    if (!MEmpty && MNoEventMode &&
+        Handler.getType() == CGType::CodeplayHostTask) {
+      assert(MGraph.expired());
+      assert(MDefaultGraphDeps.LastEventPtr == nullptr);
+      // There might be some operations submitted to the queue
+      // but the LastEventPtr is not set. If we are to run a host_task,
+      // we need to insert a barrier to ensure proper synchronization.
+      Handler.depends_on(insertHelperBarrier(Handler));
+    }
+
     // This dependency is needed for the following purposes:
     //    - host tasks are handled by the runtime and cannot be implicitly
     //    synchronized by the backend.
@@ -736,6 +771,9 @@ class queue_impl {
     //    the RT but will not be passed to the backend. See getPIEvents in
     //    Command.
     if (EventToBuildDeps) {
+      // If we have last event, this means we are no longer in no-event mode.
+      assert(!MNoEventMode);
+
       // 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.
@@ -754,12 +792,10 @@ class queue_impl {
         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);
+    MEmpty = false;
+    MNoEventMode = false;
+
+    synchronizeWithExternalEvent(Handler);
 
     auto EventRet = Handler.finalize();
     EventToBuildDeps = getSyclObjImpl(EventRet);
@@ -771,6 +807,9 @@ class queue_impl {
   event finalizeHandlerOutOfOrder(HandlerType &Handler) {
     const CGType Type = getSyclObjImpl(Handler)->MCGType;
     std::lock_guard<std::mutex> Lock{MMutex};
+
+    MEmpty = false;
+
     // The following code supports barrier synchronization if host task is
     // involved in the scenario. Native barriers cannot handle host task
     // dependency so in the case where some commands were not enqueued
@@ -1006,6 +1045,14 @@ class queue_impl {
 
   const bool MIsInorder;
 
+  // Specifies whether this queue records last event. This can only
+  // be true if the queue is in-order, the command graph is not
+  // associated with the queue and there has never been any host
+  // tasks submitted to the queue.
+  std::atomic<bool> MNoEventMode = true;
+
+  bool MEmpty = true;
+
   std::vector<EventImplPtr> MStreamsServiceEvents;
   std::mutex MStreamsServiceEventsMutex;
 

sycl/source/queue.cpp

@@ -302,7 +302,7 @@ getBarrierEventForInorderQueueHelper(const detail::QueueImplPtr QueueImpl) {
   assert(!QueueImpl->hasCommandGraph() &&
          "Should not be called in on graph recording.");
 
-  sycl::detail::optional<event> LastEvent = QueueImpl->getLastEvent();
+  sycl::detail::optional<event> LastEvent = QueueImpl->getLastEvent(QueueImpl);
   if (LastEvent)
     return *LastEvent;
 
@@ -321,10 +321,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 getBarrierEventForInorderQueueHelper(impl);
   }
 
   return submit([=](handler &CGH) { CGH.ext_oneapi_barrier(); }, CodeLoc);
@@ -349,10 +346,7 @@ event queue::ext_oneapi_submit_barrier(const std::vector<event> &WaitList,
       });
   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 getBarrierEventForInorderQueueHelper(impl);
   }
 
   return submit([=](handler &CGH) { CGH.ext_oneapi_barrier(WaitList); },
@@ -432,20 +426,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) {