Merge remote-tracking branch 'intel/sycl' into esimd_bfloat16

Author: kbobrovs

sycl/include/sycl/reduction.hpp

@@ -2272,7 +2272,7 @@ void reduCGFunc(handler &CGH, KernelType KernelFunc,
 
 namespace reduction {
 namespace aux_krn {
-template <class KernelName, class Accessor> struct Multi;
+template <class KernelName, class Predicate> struct Multi;
 } // namespace aux_krn
 } // namespace reduction
 template <typename KernelName, typename KernelType, typename... Reductions,
@@ -2312,7 +2312,7 @@ size_t reduAuxCGFunc(handler &CGH, size_t NWorkItems, size_t MaxWGSize,
     auto AccReduIndices = filterSequence<Reductions...>(Predicate, ReduIndices);
     associateReduAccsWithHandler(CGH, ReduTuple, AccReduIndices);
     using Name = __sycl_reduction_kernel<reduction::aux_krn::Multi, KernelName,
-                                         decltype(OutAccsTuple)>;
+                                         decltype(Predicate)>;
     // TODO: Opportunity to parallelize across number of elements
     range<1> GlobalRange = {HasUniformWG ? NWorkItems : NWorkGroups * WGSize};
     nd_range<1> Range{GlobalRange, range<1>(WGSize)};

sycl/source/detail/event_impl.cpp

@@ -342,10 +342,16 @@ event_impl::get_info<info::event::command_execution_status>() {
   if (MState == HES_Discarded)
     return info::event_command_status::ext_oneapi_unknown;
 
-  if (!MHostEvent && MEvent) {
-    return get_event_info<info::event::command_execution_status>(
-        this->getHandleRef(), this->getPlugin());
+  if (!MHostEvent) {
+    // Command is enqueued and PiEvent is ready
+    if (MEvent)
+      return get_event_info<info::event::command_execution_status>(
+          this->getHandleRef(), this->getPlugin());
+    // Command is blocked and not enqueued, PiEvent is not assigned yet
+    else if (MCommand)
+      return sycl::info::event_command_status::submitted;
   }
+
   return MHostEvent && MState.load() != HES_Complete
              ? sycl::info::event_command_status::submitted
              : info::event_command_status::complete;

sycl/source/detail/event_impl.hpp

@@ -212,6 +212,11 @@ class event_impl {
   }
   bool needsCleanupAfterWait() { return MNeedsCleanupAfterWait; }
 
+  /// Returns worker queue for command.
+  ///
+  /// @return a reference to MWorkerQueue.
+  QueueImplPtr &getWorkerQueue() { return MWorkerQueue; };
+
   /// Checks if an event is in a fully intialized state. Default-constructed
   /// events will return true only after having initialized its native event,
   /// while other events will assume that they are fully initialized at
@@ -243,6 +248,8 @@ class event_impl {
   std::weak_ptr<queue_impl> MQueue;
   const bool MIsProfilingEnabled = false;
 
+  QueueImplPtr MWorkerQueue;
+
   /// Dependency events prepared for waiting by backend.
   std::vector<EventImplPtr> MPreparedDepsEvents;
   std::vector<EventImplPtr> MPreparedHostDepsEvents;

sycl/source/detail/queue_impl.hpp

@@ -120,9 +120,7 @@ class queue_impl {
     }
     if (!MHostQueue) {
       const QueueOrder QOrder =
-          MPropList.has_property<property::queue::in_order>()
-              ? QueueOrder::Ordered
-              : QueueOrder::OOO;
+          MIsInorder ? QueueOrder::Ordered : QueueOrder::OOO;
       MQueues.push_back(createQueue(QOrder));
     }
   }
@@ -202,6 +200,8 @@ class queue_impl {
   /// \return true if this queue has discard_events support.
   bool has_discard_events_support() const { return MHasDiscardEventsSupport; }
 
+  bool isInOrder() const { return MIsInorder; }
+
   /// Queries SYCL queue for information.
   ///
   /// The return type depends on information being queried.

sycl/source/detail/scheduler/commands.cpp

@@ -203,17 +203,34 @@ static std::string commandToName(Command::CommandType Type) {
 }
 #endif
 
-static std::vector<RT::PiEvent>
-getPiEvents(const std::vector<EventImplPtr> &EventImpls) {
+std::vector<RT::PiEvent>
+Command::getPiEvents(const std::vector<EventImplPtr> &EventImpls) const {
   std::vector<RT::PiEvent> RetPiEvents;
   for (auto &EventImpl : EventImpls) {
-    if (EventImpl->getHandleRef() != nullptr)
-      RetPiEvents.push_back(EventImpl->getHandleRef());
+    if (EventImpl->getHandleRef() == nullptr)
+      continue;
+
+    // Do not add redundant event dependencies for in-order queues.
+    // At this stage dependency is definitely pi task and need to check if
+    // current one is a host task. In this case we should not skip pi event due
+    // to different sync mechanisms for different task types on in-order queue.
+    const QueueImplPtr &WorkerQueue = getWorkerQueue();
+    if (EventImpl->getWorkerQueue() == WorkerQueue &&
+        WorkerQueue->isInOrder() && !isHostTask())
+      continue;
+
+    RetPiEvents.push_back(EventImpl->getHandleRef());
   }
 
   return RetPiEvents;
 }
 
+bool Command::isHostTask() const {
+  return (MType == CommandType::RUN_CG) /* host task has this type also */ &&
+         ((static_cast<const ExecCGCommand *>(this))->getCG().getType() ==
+          CG::CGTYPE::CodeplayHostTask);
+}
+
 static void flushCrossQueueDeps(const std::vector<EventImplPtr> &EventImpls,
                                 const QueueImplPtr &Queue) {
   for (auto &EventImpl : EventImpls) {
@@ -240,7 +257,8 @@ class DispatchHostTask {
     // sophisticated waiting mechanism to allow to utilize this thread for any
     // other available job and resume once all required events are ready.
     for (auto &PluginWithEvents : RequiredEventsPerPlugin) {
-      std::vector<RT::PiEvent> RawEvents = getPiEvents(PluginWithEvents.second);
+      std::vector<RT::PiEvent> RawEvents =
+          MThisCmd->getPiEvents(PluginWithEvents.second);
       try {
         PluginWithEvents.first->call<PiApiKind::piEventsWait>(RawEvents.size(),
                                                               RawEvents.data());
@@ -393,10 +411,12 @@ void Command::waitForEvents(QueueImplPtr Queue,
 Command::Command(CommandType Type, QueueImplPtr Queue)
     : MQueue(std::move(Queue)),
       MEvent(std::make_shared<detail::event_impl>(MQueue)),
+      MWorkerQueue(MEvent->getWorkerQueue()),
       MPreparedDepsEvents(MEvent->getPreparedDepsEvents()),
       MPreparedHostDepsEvents(MEvent->getPreparedHostDepsEvents()),
       MType(Type) {
   MSubmittedQueue = MQueue;
+  MWorkerQueue = MQueue;
   MEvent->setCommand(this);
   MEvent->setContextImpl(MQueue->getContextImplPtr());
   MEvent->setStateIncomplete();
@@ -600,12 +620,6 @@ Command *Command::processDepEvent(EventImplPtr DepEvent, const DepDesc &Dep,
 
   Command *ConnectionCmd = nullptr;
 
-  // Do not add redundant event dependencies for in-order queues.
-  if (Dep.MDepCommand && Dep.MDepCommand->getWorkerQueue() == WorkerQueue &&
-      WorkerQueue->has_property<property::queue::in_order>() &&
-      getType() != CommandType::HOST_TASK)
-    return nullptr;
-
   ContextImplPtr DepEventContext = DepEvent->getContextImpl();
   // If contexts don't match we'll connect them using host task
   if (DepEventContext != WorkerContext && !WorkerContext->is_host()) {
@@ -621,14 +635,14 @@ const ContextImplPtr &Command::getWorkerContext() const {
   return MQueue->getContextImplPtr();
 }
 
-const QueueImplPtr &Command::getWorkerQueue() const { return MQueue; }
+const QueueImplPtr &Command::getWorkerQueue() const {
+  assert(MWorkerQueue && "MWorkerQueue must not be nullptr");
+  return MWorkerQueue;
+}
 
 bool Command::producesPiEvent() const { return true; }
 
-bool Command::supportsPostEnqueueCleanup() const {
-  // Isolated commands are cleaned up separately
-  return !MUsers.empty() || !MDeps.empty();
-}
+bool Command::supportsPostEnqueueCleanup() const { return true; }
 
 Command *Command::addDep(DepDesc NewDep, std::vector<Command *> &ToCleanUp) {
   Command *ConnectionCmd = nullptr;
@@ -1298,6 +1312,9 @@ MemCpyCommand::MemCpyCommand(Requirement SrcReq,
   if (!MSrcQueue->is_host()) {
     MEvent->setContextImpl(MSrcQueue->getContextImplPtr());
   }
+
+  MWorkerQueue = MQueue->is_host() ? MSrcQueue : MQueue;
+
   emitInstrumentationDataProxy();
 }
 
@@ -1335,10 +1352,6 @@ const ContextImplPtr &MemCpyCommand::getWorkerContext() const {
   return getWorkerQueue()->getContextImplPtr();
 }
 
-const QueueImplPtr &MemCpyCommand::getWorkerQueue() const {
-  return MQueue->is_host() ? MSrcQueue : MQueue;
-}
-
 bool MemCpyCommand::producesPiEvent() const {
   // TODO remove this workaround once the batching issue is addressed in Level
   // Zero plugin.
@@ -1481,6 +1494,8 @@ MemCpyCommandHost::MemCpyCommandHost(Requirement SrcReq,
     MEvent->setContextImpl(MSrcQueue->getContextImplPtr());
   }
 
+  MWorkerQueue = MQueue->is_host() ? MSrcQueue : MQueue;
+
   emitInstrumentationDataProxy();
 }
 
@@ -1518,10 +1533,6 @@ const ContextImplPtr &MemCpyCommandHost::getWorkerContext() const {
   return getWorkerQueue()->getContextImplPtr();
 }
 
-const QueueImplPtr &MemCpyCommandHost::getWorkerQueue() const {
-  return MQueue->is_host() ? MSrcQueue : MQueue;
-}
-
 pi_int32 MemCpyCommandHost::enqueueImp() {
   const QueueImplPtr &Queue = getWorkerQueue();
   waitForPreparedHostEvents();

sycl/source/detail/scheduler/commands.hpp

@@ -199,19 +199,28 @@ class Command {
 
   /// Get the queue this command will be submitted to. Could differ from MQueue
   /// for memory copy commands.
-  virtual const QueueImplPtr &getWorkerQueue() const;
+  const QueueImplPtr &getWorkerQueue() const;
 
   /// Returns true iff the command produces a PI event on non-host devices.
   virtual bool producesPiEvent() const;
 
   /// Returns true iff this command can be freed by post enqueue cleanup.
   virtual bool supportsPostEnqueueCleanup() const;
 
+  /// Collect PI events from EventImpls and filter out some of them in case of
+  /// in order queue
+  std::vector<RT::PiEvent>
+  getPiEvents(const std::vector<EventImplPtr> &EventImpls) const;
+
+  bool isHostTask() const;
+
 protected:
   QueueImplPtr MQueue;
   QueueImplPtr MSubmittedQueue;
   EventImplPtr MEvent;
 
+  QueueImplPtr &MWorkerQueue;
+
   /// Dependency events prepared for waiting by backend.
   /// See processDepEvent for details.
   std::vector<EventImplPtr> &MPreparedDepsEvents;
@@ -252,6 +261,10 @@ class Command {
     return MPreparedHostDepsEvents;
   }
 
+  const std::vector<EventImplPtr> &getPreparedDepsEvents() const {
+    return MPreparedDepsEvents;
+  }
+
   /// Contains list of dependencies(edges)
   std::vector<DepDesc> MDeps;
   /// Contains list of commands that depend on the command.
@@ -492,7 +505,6 @@ class MemCpyCommand : public Command {
   const Requirement *getRequirement() const final { return &MDstReq; }
   void emitInstrumentationData() final;
   const ContextImplPtr &getWorkerContext() const final;
-  const QueueImplPtr &getWorkerQueue() const final;
   bool producesPiEvent() const final;
 
 private:
@@ -517,7 +529,6 @@ class MemCpyCommandHost : public Command {
   const Requirement *getRequirement() const final { return &MDstReq; }
   void emitInstrumentationData() final;
   const ContextImplPtr &getWorkerContext() const final;
-  const QueueImplPtr &getWorkerQueue() const final;
 
 private:
   pi_int32 enqueueImp() final;

sycl/source/detail/scheduler/graph_processor.cpp

@@ -58,23 +58,20 @@ bool Scheduler::GraphProcessor::enqueueCommand(
     return false;
   }
 
-  // Recursively enqueue all the dependencies first and
-  // exit immediately if any of the commands cannot be enqueued.
-  for (DepDesc &Dep : Cmd->MDeps) {
-    if (!enqueueCommand(Dep.MDepCommand, EnqueueResult, ToCleanUp, Blocking))
-      return false;
+  // Recursively enqueue all the implicit + explicit backend level dependencies
+  // first and exit immediately if any of the commands cannot be enqueued.
+  for (const EventImplPtr &Event : Cmd->getPreparedDepsEvents()) {
+    if (Command *DepCmd = static_cast<Command *>(Event->getCommand()))
+      if (!enqueueCommand(DepCmd, EnqueueResult, ToCleanUp, Blocking))
+        return false;
   }
 
-  // Asynchronous host operations (amongst dependencies of an arbitrary command)
-  // are not supported (see Command::processDepEvent method). This impacts
-  // operation of host-task feature a lot with hangs and long-runs. Hence we
-  // have this workaround here.
-  // This workaround is safe as long as the only asynchronous host operation we
-  // have is a host task.
-  // This may iterate over some of dependencies in Cmd->MDeps. Though, the
-  // enqueue operation is idempotent and the second call will result in no-op.
-  // TODO remove the workaround when proper fix for host-task dispatching is
-  // implemented.
+  // Recursively enqueue all the implicit + explicit host dependencies and
+  // exit immediately if any of the commands cannot be enqueued.
+  // Host task execution is asynchronous. In current implementation enqueue for
+  // this command will wait till host task completion by waitInternal call on
+  // MHostDepsEvents. TO FIX: implement enqueue of blocked commands on host task
+  // completion stage and eliminate this event waiting in enqueue.
   for (const EventImplPtr &Event : Cmd->getPreparedHostDepsEvents()) {
     if (Command *DepCmd = static_cast<Command *>(Event->getCommand()))
       if (!enqueueCommand(DepCmd, EnqueueResult, ToCleanUp, Blocking))

sycl/source/detail/scheduler/scheduler.cpp

@@ -157,12 +157,6 @@ EventImplPtr Scheduler::addCG(std::unique_ptr<detail::CG> CommandGroup,
         CleanUp();
         std::rethrow_exception(std::current_exception());
       }
-
-      // If there are no memory dependencies decouple and free the command.
-      // Though, dismiss ownership of native kernel command group as it's
-      // resources may be in use by backend and synchronization point here is
-      // at native kernel execution finish.
-      CleanUp();
     }
   }
   cleanupCommands(ToCleanUp);

sycl/unittests/scheduler/CMakeLists.txt

@@ -22,4 +22,5 @@ add_sycl_unittest(SchedulerTests OBJECT
     LeafLimitDiffContexts.cpp
     InOrderQueueSyncCheck.cpp
     RunOnHostIntelCG.cpp
+    EnqueueWithDependsOnDeps.cpp
 )