alternative solution to optimize memory transfers

Author: t4c1

sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp

@@ -314,6 +314,10 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
 
   bool isInterop() const;
 
+  bool isHostPointerReadOnly() const{
+    return MHostPtrReadOnly;
+  }
+
 protected:
   // An allocateMem helper that determines which host ptr to use
   void determineHostPtr(const ContextImplPtr &Context, bool InitFromUserData,

sycl/source/detail/memory_manager.cpp

@@ -296,14 +296,10 @@ void *MemoryManager::allocateHostMemory(SYCLMemObjI *MemObj, void *UserPtr,
                                         bool HostPtrReadOnly, size_t Size,
                                         const sycl::property_list &) {
   // Can return user pointer directly if it points to writable memory.
-  if (UserPtr && HostPtrReadOnly == false)
+  if (UserPtr)
     return UserPtr;
 
   void *NewMem = MemObj->allocateHostMem();
-  // Need to initialize new memory if user provides pointer to read only
-  // memory.
-  if (UserPtr && HostPtrReadOnly == true)
-    std::memcpy((char *)NewMem, (char *)UserPtr, Size);
   return NewMem;
 }
 
@@ -331,8 +327,7 @@ static RT::PiMemFlags getMemObjCreationFlags(void *UserPtr,
   RT::PiMemFlags Result =
       HostPtrReadOnly ? PI_MEM_ACCESS_READ_ONLY : PI_MEM_FLAGS_ACCESS_RW;
   if (UserPtr)
-    Result |= HostPtrReadOnly ? PI_MEM_FLAGS_HOST_PTR_COPY
-                              : PI_MEM_FLAGS_HOST_PTR_USE;
+    Result |= PI_MEM_FLAGS_HOST_PTR_USE;
   return Result;
 }
 

sycl/source/detail/scheduler/commands.cpp

@@ -833,10 +833,10 @@ const char *Command::getBlockReason() const {
 
 AllocaCommandBase::AllocaCommandBase(CommandType Type, QueueImplPtr Queue,
                                      Requirement Req,
-                                     AllocaCommandBase *LinkedAllocaCmd)
+                                     AllocaCommandBase *LinkedAllocaCmd, bool IsConst)
     : Command(Type, Queue), MLinkedAllocaCmd(LinkedAllocaCmd),
       MIsLeaderAlloca(nullptr == LinkedAllocaCmd), MRequirement(std::move(Req)),
-      MReleaseCmd(Queue, this) {
+      MReleaseCmd(Queue, this), MIsConst(IsConst)  {
   MRequirement.MAccessMode = access::mode::read_write;
   emitInstrumentationDataProxy();
 }
@@ -868,9 +868,9 @@ bool AllocaCommandBase::supportsPostEnqueueCleanup() const { return false; }
 
 AllocaCommand::AllocaCommand(QueueImplPtr Queue, Requirement Req,
                              bool InitFromUserData,
-                             AllocaCommandBase *LinkedAllocaCmd)
+                             AllocaCommandBase *LinkedAllocaCmd, bool IsConst)
     : AllocaCommandBase(CommandType::ALLOCA, std::move(Queue), std::move(Req),
-                        LinkedAllocaCmd),
+                        LinkedAllocaCmd, IsConst),
       MInitFromUserData(InitFromUserData) {
   // Node event must be created before the dependent edge is added to this node,
   // so this call must be before the addDep() call.

sycl/source/detail/scheduler/commands.hpp

@@ -364,7 +364,8 @@ class ReleaseCommand : public Command {
 class AllocaCommandBase : public Command {
 public:
   AllocaCommandBase(CommandType Type, QueueImplPtr Queue, Requirement Req,
-                    AllocaCommandBase *LinkedAllocaCmd);
+                    AllocaCommandBase *LinkedAllocaCmd,
+                    bool IsConst = false);
 
   ReleaseCommand *getReleaseCmd() { return &MReleaseCmd; }
 
@@ -394,6 +395,8 @@ class AllocaCommandBase : public Command {
   /// Indicates that the command owns memory allocation in case of connected
   /// alloca command.
   bool MIsLeaderAlloca = true;
+  // Indicates tha thte data in this allocation must not be modified
+  bool MIsConst = false;
 
 protected:
   Requirement MRequirement;
@@ -406,7 +409,8 @@ class AllocaCommand : public AllocaCommandBase {
 public:
   AllocaCommand(QueueImplPtr Queue, Requirement Req,
                 bool InitFromUserData = true,
-                AllocaCommandBase *LinkedAllocaCmd = nullptr);
+                AllocaCommandBase *LinkedAllocaCmd = nullptr,
+                bool IsConst = false);
 
   void *getMemAllocation() const final { return MMemAllocation; }
   void printDot(std::ostream &Stream) const final;

sycl/source/detail/scheduler/graph_builder.cpp

@@ -630,14 +630,16 @@ DepDesc Scheduler::GraphBuilder::findDepForRecord(Command *Cmd,
 AllocaCommandBase *
 Scheduler::GraphBuilder::findAllocaForReq(MemObjRecord *Record,
                                           const Requirement *Req,
-                                          const ContextImplPtr &Context) {
-  auto IsSuitableAlloca = [&Context, Req](AllocaCommandBase *AllocaCmd) {
+                                          const ContextImplPtr &Context,
+                                          bool allowConst) {
+  auto IsSuitableAlloca = [&Context, Req, allowConst](AllocaCommandBase *AllocaCmd) {
     bool Res = sameCtx(AllocaCmd->getQueue()->getContextImplPtr(), Context);
     if (IsSuitableSubReq(Req)) {
       const Requirement *TmpReq = AllocaCmd->getRequirement();
       Res &= AllocaCmd->getType() == Command::CommandType::ALLOCA_SUB_BUF;
       Res &= TmpReq->MOffsetInBytes == Req->MOffsetInBytes;
       Res &= TmpReq->MSYCLMemObj->getSize() == Req->MSYCLMemObj->getSize();
+      Res &= allowConst || !AllocaCmd->MIsConst;
     }
     return Res;
   };
@@ -669,7 +671,7 @@ AllocaCommandBase *Scheduler::GraphBuilder::getOrCreateAllocaForReq(
     std::vector<Command *> &ToEnqueue) {
 
   AllocaCommandBase *AllocaCmd =
-      findAllocaForReq(Record, Req, Queue->getContextImplPtr());
+      findAllocaForReq(Record, Req, Queue->getContextImplPtr(), false);
 
   if (!AllocaCmd) {
     std::vector<Command *> ToCleanUp;
@@ -722,7 +724,7 @@ AllocaCommandBase *Scheduler::GraphBuilder::getOrCreateAllocaForReq(
                 Scheduler::getInstance().getDefaultHostQueue();
             AllocaCommand *HostAllocaCmd = new AllocaCommand(
                 DefaultHostQueue, FullReq, true /* InitFromUserData */,
-                nullptr /* LinkedAllocaCmd */);
+                nullptr /* LinkedAllocaCmd */, MemObj->isHostPointerReadOnly() /* IsConst */);
             Record->MAllocaCommands.push_back(HostAllocaCmd);
             Record->MWriteLeaves.push_back(HostAllocaCmd, ToEnqueue);
             ++(HostAllocaCmd->MLeafCounter);
@@ -755,7 +757,7 @@ AllocaCommandBase *Scheduler::GraphBuilder::getOrCreateAllocaForReq(
                                  : HostUnifiedMemory;
             if (PinnedHostMemory || HostUnifiedMemoryOnNonHostDevice) {
               AllocaCommandBase *LinkedAllocaCmdCand =
-                  findAllocaForReq(Record, Req, Record->MCurContext);
+                  findAllocaForReq(Record, Req, Record->MCurContext, false);
 
               // Cannot setup link if candidate is linked already
               if (LinkedAllocaCmdCand &&

sycl/source/detail/scheduler/scheduler.hpp

@@ -618,7 +618,8 @@ class Scheduler {
     /// Searches for suitable alloca in memory record.
     AllocaCommandBase *findAllocaForReq(MemObjRecord *Record,
                                         const Requirement *Req,
-                                        const ContextImplPtr &Context);
+                                        const ContextImplPtr &Context,
+                                        bool allowConst = true);
 
     friend class Command;