diff --git a/sycl/source/detail/scheduler/graph_builder.cpp b/sycl/source/detail/scheduler/graph_builder.cpp index 62db768e40dfc..146b52e169f35 100644 --- a/sycl/source/detail/scheduler/graph_builder.cpp +++ b/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(AllocaCmdDst)->getParentAlloca(); - else - assert( - !"Inappropriate alloca command. AllocaSubBufCommand was expected."); } AllocaCommandBase *AllocaCmdSrc = diff --git a/sycl/test/basic_tests/buffer/subbuffer.cpp b/sycl/test/basic_tests/buffer/subbuffer.cpp index 6014fa716d400..8842fd2b73464 100644 --- a/sycl/test/basic_tests/buffer/subbuffer.cpp +++ b/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 buf(a, sycl::range<1>(N)); + sycl::buffer subbuf1(buf, sycl::id<1>(0), sycl::range<1>(N / 2)); + queue1.submit([&](sycl::handler &cgh) { + auto bufacc = subbuf1.get_access(cgh); + cgh.parallel_for( + sycl::range<1>(N / 2), [=](sycl::id<1> idx) { bufacc[idx[0]] = -1; }); + }); + auto host_acc = subbuf1.get_access(); + assert(host_acc[0] == -1 && host_acc[N / 2 - 1] == -1 && + "Sub buffer data loss"); + } } int main() {