[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,17 +282,21 @@ 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 (MGraph.expired() && !MDefaultGraphDeps.LastEventPtr) {
+    assert(!MHostTaskMode);
+    // 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);
   return detail::createSyclObjFromImpl<event>(MDefaultGraphDeps.LastEventPtr);
@@ -305,7 +309,6 @@ void queue_impl::addEvent(const event &Event) {
   if (Cmd != nullptr && EImpl->getHandle() == nullptr &&
       !EImpl->isDiscarded()) {
     std::weak_ptr<event_impl> EventWeakPtr{EImpl};
-    std::lock_guard<std::mutex> Lock{MMutex};
     MEventsWeak.push_back(std::move(EventWeakPtr));
   }
 }
@@ -339,9 +342,16 @@ event queue_impl::submit_impl(const detail::type_erased_cgfo_ty &CGF,
 
   HandlerImpl->MEventMode = SubmitInfo.EventMode();
 
-  auto Event = finalizeHandler(Handler, SubmitInfo.PostProcessorFunc());
+  std::unique_lock<std::mutex> Lock(MMutex, std::defer_lock);
+  auto Event = finalizeHandler(Handler, SubmitInfo.PostProcessorFunc(), Lock);
 
-  addEvent(Event);
+  if (isInOrder() && !shouldRecordLastEvent() && Streams.empty()) {
+    // NOP
+  } else {
+    addEvent(Event);
+  }
+
+  Lock.unlock();
 
   const auto &EventImpl = detail::getSyclObjImpl(Event);
   for (auto &Stream : Streams) {
@@ -395,10 +405,13 @@ event queue_impl::submit_impl(const detail::type_erased_cgfo_ty &CGF,
 
   HandlerImpl->MEventMode = SubmitInfo.EventMode();
 
-  auto Event = finalizeHandler(Handler, SubmitInfo.PostProcessorFunc());
+  std::unique_lock<std::mutex> Lock(MMutex, std::defer_lock);
 
+  auto Event = finalizeHandler(Handler, SubmitInfo.PostProcessorFunc(), Lock);
   addEvent(Event);
 
+  Lock.unlock();
+
   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 +473,13 @@ 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 && supportsDiscardingPiEvents() &&
+          !shouldRecordLastEvent()) {
         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 +505,7 @@ event queue_impl::submitMemOpHelper(const std::shared_ptr<queue_impl> &Self,
         }
       }
 
-      if (isInOrder()) {
+      if (shouldRecordLastEvent()) {
         auto &EventToStoreIn = MGraph.expired() ? MDefaultGraphDeps.LastEventPtr
                                                 : MExtGraphDeps.LastEventPtr;
         EventToStoreIn = EventImpl;
@@ -745,14 +750,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) == MHostTaskMode);
     // 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 +767,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) {
@@ -720,40 +721,35 @@ class queue_impl {
   }
 
   template <typename HandlerType = handler>
-  event finalizeHandlerInOrder(HandlerType &Handler) {
-    // Accessing and changing of an event isn't atomic operation.
-    // Hence, here is the lock for thread-safety.
-    std::lock_guard<std::mutex> Lock{MMutex};
+  event finalizeHandlerInOrder(HandlerType &Handler,
+                               std::unique_lock<std::mutex> &Lock) {
+    Lock.lock();
 
     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);
+    if (Handler.getType() == CGType::CodeplayHostTask) {
+      if (!MHostTaskMode && MGraph.expired() && !MEmpty) {
+        assert(EventToBuildDeps == nullptr);
+        // since we don't store any events, insert a barrier to ensure proper
+        // ordering with device execution
+        auto barrierEvent = insertHelperBarrier(Handler);
+        Handler.depends_on(barrierEvent);
+      }
+
+      MHostTaskMode = true;
+    }
 
+    if (EventToBuildDeps && Handler.getType() != CGType::AsyncAlloc) {
       // 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);
+      Handler.depends_on(EventToBuildDeps);
     }
 
+    MEmpty = false;
+
     // 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.
@@ -762,15 +758,22 @@ class queue_impl {
       Handler.depends_on(*ExternalEvent);
 
     auto EventRet = Handler.finalize();
-    EventToBuildDeps = getSyclObjImpl(EventRet);
+
+    if (shouldRecordLastEvent()) {
+      EventToBuildDeps = getSyclObjImpl(EventRet);
+    }
 
     return EventRet;
   }
 
   template <typename HandlerType = handler>
-  event finalizeHandlerOutOfOrder(HandlerType &Handler) {
+  event finalizeHandlerOutOfOrder(HandlerType &Handler,
+                                  std::unique_lock<std::mutex> &Lock) {
     const CGType Type = getSyclObjImpl(Handler)->MCGType;
-    std::lock_guard<std::mutex> Lock{MMutex};
+    Lock.lock();
+
+    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
@@ -807,7 +810,8 @@ class queue_impl {
   template <typename HandlerType = handler>
   event finalizeHandlerPostProcess(
       HandlerType &Handler,
-      const optional<SubmitPostProcessF> &PostProcessorFunc) {
+      const optional<SubmitPostProcessF> &PostProcessorFunc,
+      std::unique_lock<std::mutex> &Lock) {
     bool IsKernel = Handler.getType() == CGType::Kernel;
     bool KernelUsesAssert = false;
 
@@ -818,8 +822,8 @@ class queue_impl {
           ProgramManager::getInstance().kernelUsesAssert(
               Handler.MKernelName.data());
 
-    auto Event = MIsInorder ? finalizeHandlerInOrder(Handler)
-                            : finalizeHandlerOutOfOrder(Handler);
+    auto Event = MIsInorder ? finalizeHandlerInOrder(Handler, Lock)
+                            : finalizeHandlerOutOfOrder(Handler, Lock);
 
     auto &PostProcess = *PostProcessorFunc;
 
@@ -831,12 +835,13 @@ class queue_impl {
   // template is needed for proper unit testing
   template <typename HandlerType = handler>
   event finalizeHandler(HandlerType &Handler,
-                        const optional<SubmitPostProcessF> &PostProcessorFunc) {
+                        const optional<SubmitPostProcessF> &PostProcessorFunc,
+                        std::unique_lock<std::mutex> &Lock) {
     if (PostProcessorFunc) {
-      return finalizeHandlerPostProcess(Handler, PostProcessorFunc);
+      return finalizeHandlerPostProcess(Handler, PostProcessorFunc, Lock);
     } else {
-      return MIsInorder ? finalizeHandlerInOrder(Handler)
-                        : finalizeHandlerOutOfOrder(Handler);
+      return MIsInorder ? finalizeHandlerInOrder(Handler, Lock)
+                        : finalizeHandlerOutOfOrder(Handler, Lock);
     }
   }
 
@@ -1006,6 +1011,21 @@ class queue_impl {
 
   const bool MIsInorder;
 
+  // Specifies whether this queue uses host tasks. If yes, then event
+  // from all operations need to be recorded for proper synchronization.
+  bool MHostTaskMode = false;
+
+  bool shouldRecordLastEvent() const {
+    // For in-order queues we rely on UR queue ordering.
+    // We only need to keep the event if host task are used
+    // (to ensure proper ordering).
+
+    // TODO: do not record last event for graphs as well
+    return MIsInorder && (MHostTaskMode || !MGraph.expired());
+  }
+
+  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) {

sycl/test-e2e/Graph/lit.local.cfg

@@ -1,3 +1,4 @@
 # https://github.com/intel/llvm/issues/17165
 if 'windows' in config.available_features:
    config.unsupported_features += ['arch-intel_gpu_bmg_g21']
+config.unsupported_features += ['level_zero_v2_adapter']

sycl/test-e2e/InOrderEventsExt/get_last_event.cpp

@@ -36,8 +36,13 @@ 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;
+  if (LastEvent->get_info<sycl::info::event::command_execution_status>() ==
+          sycl::info::event_command_status::complete &&
+      E->get_info<sycl::info::event::command_execution_status>() !=
+          sycl::info::event_command_status::complete) {
+    std::cout << "ext_oneapi_get_last_event() returned an event that is "
+                 "complete, but the event returned by CheckFunc() is not."
+              << std::endl;
     return 1;
   }
   return 0;

sycl/test-e2e/InorderQueue/in_order_ext_oneapi_submit_barrier.cpp

@@ -1,10 +1,5 @@
 // RUN: %{build} -o %t.out
-// RUN: env SYCL_UR_TRACE=2 %{run} %t.out 2>&1 | FileCheck %s
-
-// Test to check that we don't insert unnecessary urEnqueueEventsWaitWithBarrier
-// calls if queue is in-order and wait list is empty.
-
-// CHECK-NOT: <--- urEnqueueEventsWaitWithBarrier
+// RUN: %{run} %t.out
 
 #include <condition_variable>
 
@@ -26,13 +21,11 @@ int main() {
       cgh.single_task<class kernel1>([=]() { *Res += 9; });
     });
     auto BarrierEvent1 = Q.ext_oneapi_submit_barrier();
-    assert(BarrierEvent1 == Event1);
     auto Event2 = Q.submit([&](sycl::handler &cgh) {
       cgh.single_task<class kernel2>([=]() { *Res *= 2; });
     });
 
     auto BarrierEvent2 = Q.ext_oneapi_submit_barrier();
-    assert(BarrierEvent2 == Event2);
     BarrierEvent2.wait();
 
     // Check that kernel events are completed after waiting for barrier event.
@@ -50,7 +43,6 @@ int main() {
     auto Event1 = Q.submit(
         [&](sycl::handler &CGH) { CGH.host_task([&] { *Res += 1; }); });
     auto BarrierEvent1 = Q.ext_oneapi_submit_barrier();
-    assert(Event1 == BarrierEvent1);
     auto Event2 = Q.submit([&](sycl::handler &CGH) { CGH.fill(Res, 10, 1); });
 
     Q.wait();