[SYCL] Introduction of new scheduler.

This patch introduces new scheduler which is enabled when SCHEDULER_20
macro is set(set by default). The new scheduler is based on accessor
rather on buffer, so it will support images.

Also now there are two commands that implements moving memory to
other contexts instead of one - alloca memory and memcpy. There is new
command - release memory which will be ran during the sycl::buffer or
sycl::image object destruction and releases memory instances for the
memory object.

Signed-off-by: Vlad Romanov <[email protected]>

Author: romanovvlad

sycl/CMakeLists.txt

@@ -138,8 +138,12 @@ add_library("${SYCLLibrary}" SHARED
   "${sourceRootPath}/detail/os_util.cpp"
   "${sourceRootPath}/detail/sampler_impl.cpp"
   "${sourceRootPath}/detail/scheduler/commands.cpp"
+  "${sourceRootPath}/detail/scheduler/commands2.cpp"
   "${sourceRootPath}/detail/scheduler/printers.cpp"
   "${sourceRootPath}/detail/scheduler/scheduler.cpp"
+  "${sourceRootPath}/detail/scheduler/graph_processor.cpp"
+  "${sourceRootPath}/detail/scheduler/graph_builder.cpp"
+  "${sourceRootPath}/detail/scheduler/scheduler2.cpp"
   "${sourceRootPath}/detail/util.cpp"
   "${sourceRootPath}/context.cpp"
   "${sourceRootPath}/device.cpp"

sycl/include/CL/sycl.hpp

@@ -8,6 +8,10 @@
 
 #pragma once
 
+#ifndef SCHEDULER_10
+  #define SCHEDULER_20
+#endif
+
 #include <CL/sycl/accessor.hpp>
 #include <CL/sycl/atomic.hpp>
 #include <CL/sycl/buffer.hpp>
@@ -37,6 +41,7 @@
 #include <CL/sycl/types.hpp>
 #include <CL/sycl/version.hpp>
 
+#ifndef SCHEDULER_20
 // Do not include RT only function implementations for device code as it leads
 // to problem. Should be finally fixed when we introduce library.
 #ifndef __SYCL_DEVICE_ONLY__
@@ -46,3 +51,4 @@
 #include <CL/sycl/detail/scheduler/printers.cpp>
 #include <CL/sycl/detail/scheduler/scheduler.cpp>
 #endif //__SYCL_DEVICE_ONLY__
+#endif // !SCHEDULER_20

sycl/include/CL/sycl/detail/event_impl.hpp

@@ -59,11 +59,16 @@ class event_impl {
   // with the cl_event object stored in this class
   void setContextImpl(const ContextImplPtr &Context);
 
+  void *getCommand() { return m_Command; }
+
+  void setCommand(void *Command) { m_Command = Command; }
+
 private:
   cl_event m_Event = nullptr;
   ContextImplPtr m_Context;
   bool m_OpenCLInterop = false;
   bool m_HostEvent = true;
+  void *m_Command = nullptr;
 };
 
 } // namespace detail

sycl/include/CL/sycl/detail/kernel_impl.hpp

@@ -112,6 +112,8 @@ class kernel_impl {
                                                            Device.get(), Value);
   }
 
+  cl_kernel &getHandleRef() { return ClKernel; }
+
 private:
   cl_kernel ClKernel;
   context Context;

sycl/include/CL/sycl/detail/scheduler/commands.hpp

@@ -0,0 +1,239 @@
+//==-------------- commands.hpp - SYCL standard header file ----------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <atomic>
+#include <memory>
+#include <vector>
+
+#include <CL/sycl/detail/accessor_impl.hpp>
+#include <CL/sycl/detail/cg.hpp>
+
+namespace cl {
+namespace sycl {
+namespace detail {
+
+class queue_impl;
+class event_impl;
+class context_impl;
+
+using QueueImplPtr = std::shared_ptr<detail::queue_impl>;
+using EventImplPtr = std::shared_ptr<detail::event_impl>;
+using ContextImplPtr = std::shared_ptr<detail::context_impl>;
+
+class Command;
+class AllocaCommand;
+class ReleaseCommand;
+
+// DepDesc represents dependency between two commands
+struct DepDesc {
+  DepDesc(Command *DepCommand, Requirement *Req, AllocaCommand *AllocaCmd)
+      : MDepCommand(DepCommand), MReq(Req), MAllocaCmd(AllocaCmd) {}
+
+  friend bool operator<(const DepDesc &Lhs, const DepDesc &Rhs) {
+    return std::tie(Lhs.MReq, Lhs.MDepCommand) <
+           std::tie(Rhs.MReq, Rhs.MDepCommand);
+  }
+
+  // The actual dependency command.
+  Command *MDepCommand = nullptr;
+  // Requirement for the dependency.
+  Requirement *MReq = nullptr;
+  // Allocation command for the memory object we have requirement for.
+  // Used to simplify searching for memory handle.
+  AllocaCommand *MAllocaCmd = nullptr;
+};
+
+// The Command represents some action that needs to be performed on one or more
+// memory objects. The command has vector of Depdesc objects that represent
+// dependencies of the command. It has vector of pointer to commands that depend
+// on the command. It has pointer to sycl::queue object. And has event that is
+// associated with the command.
+class Command {
+public:
+  enum CommandType {
+    RUN_CG,
+    COPY_MEMORY,
+    ALLOCA,
+    RELEASE,
+    MAP_MEM_OBJ,
+    UNMAP_MEM_OBJ
+  };
+
+  Command(CommandType Type, QueueImplPtr Queue);
+
+  void addDep(DepDesc NewDep) {
+    if (NewDep.MDepCommand)
+      MDepsEvents.push_back(NewDep.MDepCommand->getEvent());
+    MDeps.push_back(NewDep);
+  }
+
+  void addDep(EventImplPtr Event) { MDepsEvents.push_back(std::move(Event)); }
+
+  void addUser(Command *NewUser) { MUsers.push_back(NewUser); }
+
+  // Return type of the command, e.g. Allocate, MemoryCopy.
+  CommandType getType() const { return MType; }
+
+  // The method checks if the command is enqueued, call enqueueImp if not and
+  // returns CL_SUCCESS on success.
+  cl_int enqueue();
+
+  bool isFinished();
+
+  bool isEnqueued() const { return MEnqueued; }
+
+  std::shared_ptr<queue_impl> getQueue() const { return MQueue; }
+
+  std::shared_ptr<event_impl> getEvent() const { return MEvent; }
+
+protected:
+  EventImplPtr MEvent;
+  QueueImplPtr MQueue;
+  std::vector<EventImplPtr> MDepsEvents;
+
+  std::vector<cl_event> prepareEvents(ContextImplPtr Context);
+
+  // Private interface. Derived classes should implement this method.
+  virtual cl_int enqueueImp() = 0;
+
+public:
+  std::vector<DepDesc> MDeps;
+  std::vector<Command *> MUsers;
+
+private:
+  CommandType MType;
+  std::atomic<bool> MEnqueued;
+};
+
+// The command enqueues release instance of memory allocated on Host or
+// underlying framework.
+class ReleaseCommand : public Command {
+public:
+  ReleaseCommand(QueueImplPtr Queue, AllocaCommand *AllocaCmd)
+      : Command(CommandType::RELEASE, std::move(Queue)), MAllocaCmd(AllocaCmd) {
+  }
+private:
+  cl_int enqueueImp() override;
+
+  AllocaCommand *MAllocaCmd = nullptr;
+};
+
+// The command enqueues allocation of instance of memory object on Host or
+// underlying framework.
+class AllocaCommand : public Command {
+public:
+  AllocaCommand(QueueImplPtr Queue, Requirement Req,
+                bool InitFromUserData = true)
+      : Command(CommandType::ALLOCA, Queue), MReleaseCmd(Queue, this),
+        MInitFromUserData(InitFromUserData), MReq(std::move(Req)) {
+    addDep(DepDesc(nullptr, &MReq, this));
+  }
+  ReleaseCommand *getReleaseCmd() { return &MReleaseCmd; }
+
+  SYCLMemObjT *getSYCLMemObj() const { return MReq.MSYCLMemObj; }
+
+  void *getMemAllocation() const { return MMemAllocation; }
+
+  Requirement *getAllocationReq() { return &MReq; }
+
+private:
+  cl_int enqueueImp() override;
+
+  ReleaseCommand MReleaseCmd;
+  void *MMemAllocation = nullptr;
+  bool MInitFromUserData = false;
+  Requirement MReq;
+};
+
+class MapMemObject : public Command {
+public:
+  MapMemObject(Requirement SrcReq, AllocaCommand *SrcAlloca,
+               Requirement *DstAcc, QueueImplPtr Queue);
+
+  Requirement MSrcReq;
+  AllocaCommand *MSrcAlloca = nullptr;
+  Requirement *MDstAcc = nullptr;
+  Requirement MDstReq;
+
+private:
+  cl_int enqueueImp() override;
+};
+
+class UnMapMemObject : public Command {
+public:
+  UnMapMemObject(Requirement SrcReq, AllocaCommand *SrcAlloca,
+                 Requirement *DstAcc, QueueImplPtr Queue);
+
+private:
+  cl_int enqueueImp() override;
+
+  Requirement MSrcReq;
+  AllocaCommand *MSrcAlloca = nullptr;
+  Requirement *MDstAcc = nullptr;
+};
+
+// The command enqueues memory copy between two instances of memory object.
+class MemCpyCommand : public Command {
+public:
+  MemCpyCommand(Requirement SrcReq, AllocaCommand *SrcAlloca,
+                Requirement DstReq, AllocaCommand *DstAlloca,
+                QueueImplPtr SrcQueue, QueueImplPtr DstQueue);
+
+  QueueImplPtr MSrcQueue;
+  Requirement MSrcReq;
+  AllocaCommand *MSrcAlloca = nullptr;
+  Requirement MDstReq;
+  AllocaCommand *MDstAlloca = nullptr;
+  Requirement *MAccToUpdate = nullptr;
+
+  void setAccessorToUpdate(Requirement *AccToUpdate) {
+    MAccToUpdate = AccToUpdate;
+  }
+
+private:
+  cl_int enqueueImp() override;
+};
+
+// The command enqueues memory copy between two instances of memory object.
+class MemCpyCommandHost : public Command {
+public:
+  MemCpyCommandHost(Requirement SrcReq, AllocaCommand *SrcAlloca,
+                    Requirement *DstAcc, QueueImplPtr SrcQueue,
+                    QueueImplPtr DstQueue);
+
+  QueueImplPtr MSrcQueue;
+  Requirement MSrcReq;
+  AllocaCommand *MSrcAlloca = nullptr;
+  Requirement MDstReq;
+  Requirement *MDstAcc = nullptr;
+
+private:
+  cl_int enqueueImp() override;
+};
+
+// The command enqueues execution of kernel or explicit memory operation.
+class ExecCGCommand : public Command {
+public:
+  ExecCGCommand(std::unique_ptr<detail::CG> CommandGroup, QueueImplPtr Queue)
+      : Command(CommandType::RUN_CG, std::move(Queue)),
+        MCommandGroup(std::move(CommandGroup)) {}
+
+private:
+  // Implementation of enqueueing of ExecCGCommand.
+  cl_int enqueueImp() override;
+
+  AllocaCommand *getAllocaForReq(Requirement *Req);
+
+  std::unique_ptr<detail::CG> MCommandGroup;
+};
+
+} // namespace detail
+} // namespace sycl
+} // namespace cl

sycl/include/CL/sycl/detail/scheduler/scheduler.h

@@ -6,6 +6,12 @@
 //
 //===----------------------------------------------------------------------===//
 
+#ifdef SCHEDULER_20
+
+#include <CL/sycl/detail/scheduler/scheduler.hpp>
+
+#else
+
 #pragma once
 
 #include <CL/sycl/context.hpp>
@@ -240,3 +246,4 @@ class Scheduler {
 } // namespace simple_scheduler
 } // namespace sycl
 } // namespace cl
+#endif // SCHEDULER_20