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

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.

Author: igchor

sycl/include/sycl/handler.hpp

@@ -468,11 +468,11 @@ class __SYCL_EXPORT handler {
   detail::CGType getType() const;
 
   void throwIfActionIsCreated() {
-    if (detail::CGType::None != getType())
-      throw sycl::exception(make_error_code(errc::runtime),
-                            "Attempt to set multiple actions for the "
-                            "command group. Command group must consist of "
-                            "a single kernel or explicit memory operation.");
+    //if (detail::CGType::None != getType())
+    //  throw sycl::exception(make_error_code(errc::runtime),
+    //                        "Attempt to set multiple actions for the "
+    //                        "command group. Command group must consist of "
+    //                        "a single kernel or explicit memory operation.");
   }
 
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES

sycl/source/detail/queue_impl.cpp

@@ -283,22 +283,6 @@ 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.
-  if (std::optional<event> ExternalEvent = MInOrderExternalEvent.read())
-    return ExternalEvent;
-
-  std::lock_guard<std::mutex> Lock{MMutex};
-  if (MGraph.expired() && !MDefaultGraphDeps.LastEventPtr)
-    return std::nullopt;
-  if (MDiscardEvents)
-    return createDiscardedEvent();
-  if (!MGraph.expired() && MExtGraphDeps.LastEventPtr)
-    return detail::createSyclObjFromImpl<event>(MExtGraphDeps.LastEventPtr);
-  return detail::createSyclObjFromImpl<event>(MDefaultGraphDeps.LastEventPtr);
-}
-
 void queue_impl::addEvent(const event &Event) {
   const EventImplPtr &EImpl = getSyclObjImpl(Event);
   assert(EImpl && "Event implementation is missing");

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>() &&
@@ -785,39 +783,28 @@ class queue_impl {
     return ResEvent;
   }
 
+  template <typename HandlerType = handler>
+  void insertHelperBarrierNoEvent(const HandlerType &Handler) {
+    getAdapter()->call<UrApiKind::urEnqueueEventsWaitWithBarrier>(
+        Handler.MQueue->getHandleRef(), 0, nullptr, nullptr);
+  }
+
   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};
 
-    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 (LastHostTaskEvent) {
+      insertHelperBarrierNoEvent(Handler);
+      LastHostTaskEvent = nullptr;
     }
 
     // If there is an external event set, add it as a dependency and clear it.
@@ -828,7 +815,10 @@ class queue_impl {
       Handler.depends_on(*ExternalEvent);
 
     auto EventRet = Handler.finalize();
-    EventToBuildDeps = getSyclObjImpl(EventRet);
+
+    if (getSyclObjImpl(EventRet)->isHost()) {
+      LastHostTaskEvent = getSyclObjImpl(EventRet);
+    }
 
     return EventRet;
   }
@@ -1012,6 +1002,9 @@ class queue_impl {
   /// need to emulate it with multiple native in-order queues.
   bool MEmulateOOO = false;
 
+  // TODO: this is for in-order queue only. Move it to separate struct.
+  EventImplPtr LastHostTaskEvent = nullptr;
+
   // Access should be guarded with MMutex
   struct DependencyTrackingItems {
     // This event is employed for enhanced dependency tracking with in-order

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,8 @@ 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;
+  // We insert a marker to represent an event at end.
+  return detail::createSyclObjFromImpl<event>(impl->insertMarkerEvent(impl));
 }
 
 void queue::ext_oneapi_set_external_event(const event &external_event) {