[SYCL] Fixed sub-buffer memory allocation update (#1486)

In some cases parent`s memory allocation might change (e.g., after
map/unmap operations). If parent`s memory allocation changes, sub-buffer
memory allocation should be changed as well.

Signed-off-by: Ivan Karachun <[email protected]>

Author: Ivan Karachun

sycl/source/detail/scheduler/commands.cpp

@@ -694,6 +694,18 @@ void AllocaSubBufCommand::emitInstrumentationData() {
 #endif
 }
 
+void *AllocaSubBufCommand::getMemAllocation() const {
+  // In some cases parent`s memory allocation might change (e.g., after
+  // map/unmap operations). If parent`s memory allocation changes, sub-buffer
+  // memory allocation should be changed as well.
+  if (MQueue->is_host()) {
+    return static_cast<void *>(
+        static_cast<char *>(MParentAlloca->getMemAllocation()) +
+        MRequirement.MOffsetInBytes);
+  }
+  return MMemAllocation;
+}
+
 cl_int AllocaSubBufCommand::enqueueImp() {
   std::vector<EventImplPtr> EventImpls =
       Command::prepareEvents(detail::getSyclObjImpl(MQueue->get_context()));

sycl/source/detail/scheduler/commands.hpp

@@ -264,7 +264,7 @@ class AllocaCommandBase : public Command {
 
   SYCLMemObjI *getSYCLMemObj() const { return MRequirement.MSYCLMemObj; }
 
-  void *getMemAllocation() const { return MMemAllocation; }
+  virtual void *getMemAllocation() const = 0;
 
   const Requirement *getRequirement() const final { return &MRequirement; }
 
@@ -298,6 +298,7 @@ class AllocaCommand : public AllocaCommandBase {
                 bool InitFromUserData = true,
                 AllocaCommandBase *LinkedAllocaCmd = nullptr);
 
+  void *getMemAllocation() const final { return MMemAllocation; }
   void printDot(std::ostream &Stream) const final;
   void emitInstrumentationData();
 
@@ -314,6 +315,7 @@ class AllocaSubBufCommand : public AllocaCommandBase {
   AllocaSubBufCommand(QueueImplPtr Queue, Requirement Req,
                       AllocaCommandBase *ParentAlloca);
 
+  void *getMemAllocation() const final;
   void printDot(std::ostream &Stream) const final;
   AllocaCommandBase *getParentAlloca() { return MParentAlloca; }
   void emitInstrumentationData();

sycl/source/detail/scheduler/graph_builder.cpp

@@ -533,6 +533,7 @@ AllocaCommandBase *Scheduler::GraphBuilder::findAllocaForReq(
     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();
     }

sycl/test/basic_tests/buffer/subbuffer.cpp

@@ -280,7 +280,7 @@ void checkMultipleContexts() {
   {
     sycl::queue queue1;
     sycl::buffer<int, 1> buf(a, sycl::range<1>(N));
-    sycl::buffer<int, 1> subbuf1(buf, sycl::id<1>(0), sycl::range<1>(N / 2));
+    sycl::buffer<int, 1> subbuf1(buf, sycl::id<1>(N / 2), sycl::range<1>(N / 2));
     queue1.submit([&](sycl::handler &cgh) {
       auto bufacc = subbuf1.get_access<sycl::access::mode::read_write>(cgh);
       cgh.parallel_for<class sub_buffer_3>(