[SYCL] Fixed sub-buffer alloca search (#1385)

If found alloca command is not sub-buffer alloca, then
it's parent alloca which has same context

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

Author: Ivan Karachun

sycl/source/detail/scheduler/graph_builder.cpp

@@ -240,9 +240,6 @@ Command *Scheduler::GraphBuilder::insertMemoryMove(MemObjRecord *Record,
     if (AllocaCmdDst->getType() == Command::CommandType::ALLOCA_SUB_BUF)
       AllocaCmdDst =
           static_cast<AllocaSubBufCommand *>(AllocaCmdDst)->getParentAlloca();
-    else
-      assert(
-          !"Inappropriate alloca command. AllocaSubBufCommand was expected.");
   }
 
   AllocaCommandBase *AllocaCmdSrc =

sycl/test/basic_tests/buffer/subbuffer.cpp

@@ -275,6 +275,20 @@ void checkMultipleContexts() {
     });
   }
   assert(a[N / 2 - 1] == 1 && a[N / 2] == 2 && "Sub buffer data loss");
+
+  {
+    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));
+    queue1.submit([&](sycl::handler &cgh) {
+      auto bufacc = subbuf1.get_access<sycl::access::mode::read_write>(cgh);
+      cgh.parallel_for<class sub_buffer_3>(
+          sycl::range<1>(N / 2), [=](sycl::id<1> idx) { bufacc[idx[0]] = -1; });
+    });
+    auto host_acc = subbuf1.get_access<sycl::access::mode::read>();
+    assert(host_acc[0] == -1 && host_acc[N / 2 - 1] == -1 &&
+           "Sub buffer data loss");
+  }
 }
 
 int main() {