[SYCL][Scheduler] Rework of host accessor and host allocation (#724)

* [SYCL][Scheduler] Rework of host accessor and host allocation

Background for use_host_ptr property:

The SYCL specification defines use_host_ptr property which can be passed
into the buffer or image constructors. The SYCL specification says the
following:

The use_host_ptr property adds the requirement that the SYCL runtime
must not allocate any memory for the SYCL buffer and instead uses the
provided host pointer directly. This prevents the SYCL runtime from
allocating additional temporary storage on the host.

Current implementation violates this rule in some scenarios, such as
moving the latest data of the buffer from one OCL device context to
another and on update_host type of explicit memory operation.

Background for host accessor:

Current implementation of host accessor requires that the underlying
backend supports async/delayed execution: SYCL RT calls Unmap or Write
commands that are initially blocked by an user event. Because of that,
for instance, host accessors work incorrectly if the latest memory are
located on the host device. Also, user events are usually not very well
supported, there is "UseExclusiveQueue" workaround because of that.

Solution:

The patch eliminates the need to have more then one host allocation by
introducing concept of linked alloca commands. One device and one host
alloca commands can be linked together so they share* the same host
memory. Such commands have "active" state which indicates that it's
valid to work with corresponding memory allocation. Only one alloca
command from a pair can be active at the same time. Memory transition
between them is performed using map/unmap commands.

* share from the SYCL point of view, because even if SYCL RT asks
underlying backends to reuse host pointer they are free to allocate
additional memory on their own.

In order to resolve the limitations of current implementation of the
host accessor the patch updates addHostAccessor API to emit map or
read operations along with an empty command which is in blocked
state(cannot be enqueued). An empty command is unblocked inside in the
new releaseHostAccessor API which is called from host accessor
destructor.

As bonuses/side effects the patch:
1. Allows having multiple host accessor with read access mode at the
   same time in all cases.
2. Adds environment variable SYCL_THROW_ON_BLOCK which if set makes SYCL
   RT throw an exception on attempt to wait for a blocked command.
3. Makes so the latest state of memory object is still on the host after
   host accessor destructor in all cases, so when multiple host
   accessors are constructed in a row no map/read operations are
   performed starting from the second host accessor. Before the patch
   memory was moved back to the source(for example, to OpenCL memory
   allocation) so SYCL RT had to make map/read each time.
4. Makes Scheduler considering host and different instances of the host
   device to be in the same context when it looks for a suitable
   allocation.

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

Author: romanovvlad

sycl/include/CL/sycl/detail/accessor_impl.hpp

@@ -20,6 +20,8 @@ namespace cl {
 namespace sycl {
 namespace detail {
 
+class Command;
+
 // The class describes a requirement to access a SYCL memory object such as
 // sycl::buffer and sycl::image. For example, each accessor used in a kernel,
 // except one with access target "local", adds such requirement for the command
@@ -70,10 +72,8 @@ class AccessorImplHost {
         MElemSize(ElemSize), MOffsetInBytes(OffsetInBytes),
         MIsSubBuffer(IsSubBuffer) {}
 
-  ~AccessorImplHost() {
-    if (BlockingEvent)
-      BlockingEvent->setComplete();
-  }
+  ~AccessorImplHost();
+
   AccessorImplHost(const AccessorImplHost &Other)
       : MOffset(Other.MOffset), MAccessRange(Other.MAccessRange),
         MMemoryRange(Other.MMemoryRange), MAccessMode(Other.MAccessMode),
@@ -97,7 +97,7 @@ class AccessorImplHost {
 
   void *MData = nullptr;
 
-  EventImplPtr BlockingEvent;
+  Command *MBlockedCmd = nullptr;
 };
 
 using AccessorImplPtr = std::shared_ptr<AccessorImplHost>;

sycl/include/CL/sycl/detail/buffer_impl.hpp

@@ -143,9 +143,17 @@ class buffer_impl final : public SYCLMemObjT<AllocatorT> {
   }
 
   void *allocateMem(ContextImplPtr Context, bool InitFromUserData,
-                    RT::PiEvent &OutEventToWait) override {
+                    void *HostPtr, RT::PiEvent &OutEventToWait) override {
 
-    void *UserPtr = InitFromUserData ? BaseT::getUserPtr() : nullptr;
+    assert(!(InitFromUserData && HostPtr) &&
+           "Cannot init from user data and reuse host ptr provided "
+           "simultaneously");
+
+    void *UserPtr = InitFromUserData ? BaseT::getUserPtr() : HostPtr;
+
+    assert(!(nullptr == UserPtr && BaseT::useHostPtr() && Context->is_host()) &&
+           "Internal error. Allocating memory on the host "
+           "while having use_host_ptr property");
 
     return MemoryManager::allocateMemBuffer(
         std::move(Context), this, UserPtr, BaseT::MHostPtrReadOnly,

sycl/include/CL/sycl/detail/image_impl.hpp

@@ -280,8 +280,13 @@ class image_impl final : public SYCLMemObjT<AllocatorT> {
   size_t get_count() const { return MRange.size(); }
 
   void *allocateMem(ContextImplPtr Context, bool InitFromUserData,
-                    RT::PiEvent &OutEventToWait) override {
-    void *UserPtr = InitFromUserData ? BaseT::getUserPtr() : nullptr;
+                    void *HostPtr, RT::PiEvent &OutEventToWait) override {
+
+    assert(!(InitFromUserData && HostPtr) &&
+           "Cannot init from user data and reuse host ptr provided "
+           "simultaneously");
+
+    void *UserPtr = InitFromUserData ? BaseT::getUserPtr() : HostPtr;
 
     RT::PiMemImageDesc Desc = getImageDesc(UserPtr != nullptr);
     assert(checkImageDesc(Desc, Context, UserPtr) &&

sycl/include/CL/sycl/detail/memory_manager.hpp

@@ -42,7 +42,8 @@ class MemoryManager {
   // The following method allocates memory allocation of memory object.
   // Depending on the context it allocates memory on host or on device.
   static void *allocate(ContextImplPtr TargetContext, SYCLMemObjI *MemObj,
-                        bool InitFromUserData, std::vector<RT::PiEvent> DepEvents,
+                        bool InitFromUserData, void *HostPtr,
+                        std::vector<RT::PiEvent> DepEvents,
                         RT::PiEvent &OutEvent);
 
   // The following method creates OpenCL sub buffer for specified

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

@@ -155,6 +155,8 @@ class Command {
   bool MIsBlockable = false;
   // Indicates whether the command is blocked from enqueueing
   std::atomic<bool> MCanEnqueue;
+
+  const char *MBlockReason = "Unknown";
 };
 
 // The command does nothing during enqueue. The task can be used to implement
@@ -193,8 +195,10 @@ class ReleaseCommand : public Command {
 
 class AllocaCommandBase : public Command {
 public:
-  AllocaCommandBase(CommandType Type, QueueImplPtr Queue, Requirement Req)
-      : Command(Type, Queue), MReleaseCmd(Queue, this),
+  AllocaCommandBase(CommandType Type, QueueImplPtr Queue, Requirement Req,
+                    AllocaCommandBase *LinkedAllocaCmd)
+      : Command(Type, Queue), MLinkedAllocaCmd(LinkedAllocaCmd),
+        MIsLeaderAlloca(nullptr == LinkedAllocaCmd), MReleaseCmd(Queue, this),
         MRequirement(std::move(Req)) {
     MRequirement.MAccessMode = access::mode::read_write;
   }
@@ -207,20 +211,35 @@ class AllocaCommandBase : public Command {
 
   const Requirement *getRequirement() const final { return &MRequirement; }
 
+  void *MMemAllocation = nullptr;
+
+  // Alloca command linked with current command.
+  // Device and host alloca commands can be linked, so they may share the same
+  // memory. Only one allocation from a pair can be accessed at a time. Alloca
+  // commands associated with such allocation is "active". In order to switch
+  // "active" status between alloca commands map/unmap operations are used.
+  AllocaCommandBase *MLinkedAllocaCmd = nullptr;
+  // Indicates that current alloca is active one.
+  bool MIsActive = true;
+
+  // Indicates that the command owns memory allocation in case of connected
+  // alloca command
+  bool MIsLeaderAlloca = true;
+
 protected:
   ReleaseCommand MReleaseCmd;
   Requirement MRequirement;
-  void *MMemAllocation = nullptr;
 };
 
 // The command enqueues allocation of instance of memory object on Host or
 // underlying framework.
 class AllocaCommand : public AllocaCommandBase {
 public:
   AllocaCommand(QueueImplPtr Queue, Requirement Req,
-                bool InitFromUserData = true)
-      : AllocaCommandBase(CommandType::ALLOCA, std::move(Queue),
-                          std::move(Req)),
+                bool InitFromUserData = true,
+                AllocaCommandBase *LinkedAllocaCmd = nullptr)
+      : AllocaCommandBase(CommandType::ALLOCA, std::move(Queue), std::move(Req),
+                          LinkedAllocaCmd),
         MInitFromUserData(InitFromUserData) {
     addDep(DepDesc(nullptr, getRequirement(), this));
   }
@@ -240,7 +259,8 @@ class AllocaSubBufCommand : public AllocaCommandBase {
   AllocaSubBufCommand(QueueImplPtr Queue, Requirement Req,
                       AllocaCommandBase *ParentAlloca)
       : AllocaCommandBase(CommandType::ALLOCA_SUB_BUF, std::move(Queue),
-                          std::move(Req)),
+                          std::move(Req),
+                          /*LinkedAllocaCmd*/ nullptr),
         MParentAlloca(ParentAlloca) {
     addDep(DepDesc(MParentAlloca, getRequirement(), MParentAlloca));
   }
@@ -251,7 +271,7 @@ class AllocaSubBufCommand : public AllocaCommandBase {
 private:
   cl_int enqueueImp() final;
 
-  AllocaCommandBase *MParentAlloca;
+  AllocaCommandBase *MParentAlloca = nullptr;
 };
 
 class MapMemObject : public Command {
@@ -295,10 +315,6 @@ class MemCpyCommand : public Command {
                 QueueImplPtr SrcQueue, QueueImplPtr DstQueue,
                 bool UseExclusiveQueue = false);
 
-  void setAccessorToUpdate(Requirement *AccToUpdate) {
-    MAccToUpdate = AccToUpdate;
-  }
-
   void printDot(std::ostream &Stream) const final;
   const Requirement *getRequirement() const final { return &MDstReq; }
 
@@ -310,7 +326,6 @@ class MemCpyCommand : public Command {
   AllocaCommandBase *MSrcAllocaCmd = nullptr;
   Requirement MDstReq;
   AllocaCommandBase *MDstAllocaCmd = nullptr;
-  Requirement *MAccToUpdate = nullptr;
 };
 
 // The command enqueues memory copy between two instances of memory object.

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

@@ -71,8 +71,17 @@ class Scheduler {
   // sycl::image destructors.
   void removeMemoryObject(detail::SYCLMemObjI *MemObj);
 
+  // Creates nodes in the graph, that update Req with the pointer to the host
+  // memory which contains the latest data of the memory object. New operations
+  // with the same memory object that have side effects are blocked until
+  // releaseHostAccessor is called.
+  // Returns an event which indicates when these nodes are completed and host
+  // accessor is ready for using.
   EventImplPtr addHostAccessor(Requirement *Req);
 
+  // Unblocks operations with the memory object.
+  void releaseHostAccessor(Requirement *Req);
+
   // Returns an instance of the scheduler object.
   static Scheduler &getInstance();
 
@@ -101,7 +110,7 @@ class Scheduler {
                              QueueImplPtr HostQueue);
 
     Command *addCopyBack(Requirement *Req);
-    Command *addHostAccessor(Requirement *Req, EventImplPtr &RetEvent);
+    Command *addHostAccessor(Requirement *Req);
 
     // [Provisional] Optimizes the whole graph.
     void optimize();
@@ -142,18 +151,20 @@ class Scheduler {
     std::vector<SYCLMemObjI *> MMemObjs;
 
   private:
-    // The method inserts memory copy operation from the context where the
-    // memory current lives to the context bound to Queue.
-    MemCpyCommand *insertMemCpyCmd(MemObjRecord *Record, Requirement *Req,
-                                   const QueueImplPtr &Queue,
-                                   bool UseExclusiveQueue = false);
+    // The method inserts required command to make so the latest state for the
+    // memory object Record refers to resides in the context which is bound to
+    // the Queue. Can insert copy/map/unmap operations depending on the source
+    // and destination.
+    Command *insertMemoryMove(MemObjRecord *Record, Requirement *Req,
+                             const QueueImplPtr &Queue,
+                             bool UseExclusiveQueue = false);
 
     UpdateHostRequirementCommand *
     insertUpdateHostReqCmd(MemObjRecord *Record, Requirement *Req,
                            const QueueImplPtr &Queue);
 
     std::set<Command *> findDepsForReq(MemObjRecord *Record, Requirement *Req,
-                                       QueueImplPtr Context);
+                                       const ContextImplPtr &Context);
 
     // Searches for suitable alloca in memory record.
     AllocaCommandBase *findAllocaForReq(MemObjRecord *Record, Requirement *Req,

sycl/include/CL/sycl/detail/sycl_mem_obj_i.hpp

@@ -38,11 +38,12 @@ class SYCLMemObjI {
   // point to event that should be waited before using the memory.
   // InitFromUserData indicates that the returned memory should be intialized
   // with the data provided by user(if any). Usually it should happen on the
-  // first allocation of memory for the buffer.
+  // first allocation of memory for the memory object.
+  // Non null HostPtr requires allocation to be made with USE_HOST_PTR property.
   // Method returns a pointer to host allocation if Context is host one and
   // cl_mem obect if not.
   virtual void *allocateMem(ContextImplPtr Context, bool InitFromUserData,
-                            RT::PiEvent &InteropEvent) = 0;
+                            void *HostPtr, RT::PiEvent &InteropEvent) = 0;
 
   // Should be used for memory object created without use_host_ptr property.
   virtual void *allocateHostMem() = 0;

sycl/source/CMakeLists.txt

@@ -34,6 +34,7 @@ endfunction(add_sycl_rt_library)
 
 set(SYCL_SOURCES
     "${sycl_inc_dir}/CL/sycl.hpp"
+    "detail/accessor_impl.cpp"
     "detail/builtins_common.cpp"
     "detail/builtins_geometric.cpp"
     "detail/builtins_integer.cpp"

sycl/source/detail/accessor_impl.cpp

@@ -0,0 +1,23 @@
+//==---------------- accessor_impl.cpp - SYCL standard source 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
+//
+//===----------------------------------------------------------------------===//
+
+#include <CL/sycl/detail/accessor_impl.hpp>
+#include <CL/sycl/detail/scheduler/scheduler.hpp>
+
+namespace cl {
+namespace sycl {
+namespace detail {
+
+AccessorImplHost::~AccessorImplHost() {
+  if (MBlockedCmd)
+    detail::Scheduler::getInstance().releaseHostAccessor(this);
+}
+}
+}
+}
+

sycl/source/detail/event_impl.cpp

@@ -36,10 +36,6 @@ event_impl::~event_impl() {
   }
 }
 
-void event_impl::setComplete() {
-  PI_CALL(RT::piEventSetStatus, m_Event, CL_COMPLETE);
-}
-
 void event_impl::waitInternal() const {
   if (!m_HostEvent) {
     PI_CALL(RT::piEventsWait, 1, &m_Event);

sycl/source/detail/memory_manager.cpp

@@ -54,15 +54,16 @@ void MemoryManager::releaseMemObj(ContextImplPtr TargetContext,
 }
 
 void *MemoryManager::allocate(ContextImplPtr TargetContext, SYCLMemObjI *MemObj,
-                              bool InitFromUserData,
+                              bool InitFromUserData, void *HostPtr,
                               std::vector<RT::PiEvent> DepEvents,
                               RT::PiEvent &OutEvent) {
   // There is no async API for memory allocation. Explicitly wait for all
   // dependency events and return empty event.
   waitForEvents(DepEvents);
   OutEvent = nullptr;
 
-  return MemObj->allocateMem(TargetContext, InitFromUserData, OutEvent);
+  return MemObj->allocateMem(TargetContext, InitFromUserData, HostPtr,
+                             OutEvent);
 }
 
 void *MemoryManager::allocateHostMemory(SYCLMemObjI *MemObj, void *UserPtr,
@@ -420,7 +421,7 @@ void *MemoryManager::map(SYCLMemObjI *SYCLMemObj, void *Mem, QueueImplPtr Queue,
                          sycl::id<3> AccessOffset, unsigned int ElementSize,
                          std::vector<RT::PiEvent> DepEvents,
                          RT::PiEvent &OutEvent) {
-  if (Queue->is_host() || Dim != 1) {
+  if (Queue->is_host()) {
     assert(!"Not supported configuration of map requested");
     throw runtime_error("Not supported configuration of map requested");
   }
@@ -447,10 +448,15 @@ void *MemoryManager::map(SYCLMemObjI *SYCLMemObj, void *Mem, QueueImplPtr Queue,
   AccessOffset[0] *= ElementSize;
   AccessRange[0] *= ElementSize;
 
-  void *MappedPtr;
+  // TODO: Handle offset
+  assert(AccessOffset[0] == 0 && "Handle offset");
+
+  void *MappedPtr = nullptr;
+  const size_t BytesToMap = AccessRange[0] * AccessRange[1] * AccessRange[2];
+
   PI_CALL(RT::piEnqueueMemBufferMap, Queue->getHandleRef(),
           pi::cast<RT::PiMem>(Mem), CL_FALSE, Flags, AccessOffset[0],
-          AccessRange[0], DepEvents.size(),
+          BytesToMap, DepEvents.size(),
           DepEvents.empty() ? nullptr : &DepEvents[0], &OutEvent, &MappedPtr);
   return MappedPtr;
 }