[SYCL] Test for fix of linked alloca's deps (#1470)

The tests provided here are for the fix introduced with #1471.

Use case fixed by the fix. Imagine us having two device queues (`Q1` and `Q2`) and a single buffer `B`. We initialize the buffer with host accessor. Also lets have two kernels: kernel `K1` executed at `Q1`; and kernel `K2` which executes at `Q2`. `K1` writes to `B`, `K2` reads from `B`. After submitting `K1` to `Q1` there will also be an `AllocaCommand` `A1` which allocates buffer on device. `K1` depends on `A1` via memory object `B`. After submitting `K2` to `Q2` there will be another `AllocaCommand` `A2` (for queue `Q2` and its device). `K2` will depend on `A2`. `A2`, however, should depend on both `A1` and `K1`. `A2->K1` dependency eliminates data race.

Signed-off-by: Sergey Kanaev <[email protected]>

Author: s-kanaev

sycl/source/detail/scheduler/scheduler.hpp

@@ -166,6 +166,9 @@
 /// clReleaseContext(ContextCPU);
 /// \endcode
 
+// For testing purposes
+class MockScheduler;
+
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
@@ -576,6 +579,8 @@ class Scheduler {
     friend class Command;
 
   private:
+    friend class ::MockScheduler;
+
     /// Searches for suitable alloca in memory record.
     ///
     /// If none found, creates new one.

sycl/unittests/scheduler/CMakeLists.txt

@@ -6,5 +6,6 @@ add_sycl_unittest(SchedulerTests OBJECT
     MemObjCommandCleanup.cpp
     CommandsWaitForEvents.cpp
     WaitAfterCleanup.cpp
+    LinkedAllocaDependencies.cpp
     utils.cpp
 )

sycl/unittests/scheduler/LinkedAllocaDependencies.cpp

@@ -0,0 +1,96 @@
+//==------ LinkedAllocaDependencies.cpp --- Scheduler unit tests -----------==//
+//
+// 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 "SchedulerTest.hpp"
+#include "SchedulerTestUtils.hpp"
+
+using namespace cl::sycl;
+
+class MemObjMock : public cl::sycl::detail::SYCLMemObjI {
+public:
+  using ContextImplPtr = std::shared_ptr<cl::sycl::detail::context_impl>;
+
+  MemObjMock(const std::shared_ptr<cl::sycl::detail::MemObjRecord> &Record)
+      : SYCLMemObjI() {
+    MRecord = Record;
+  }
+
+  ~MemObjMock() = default;
+
+  MemObjType getType() const override { return MemObjType::BUFFER; }
+
+  void *allocateMem(ContextImplPtr, bool, void *,
+                    cl::sycl::detail::pi::PiEvent &) {
+    return nullptr;
+  }
+
+  void *allocateHostMem() { return nullptr; }
+  void releaseMem(ContextImplPtr, void *) {}
+  void releaseHostMem(void *) {}
+  size_t getSize() const override { return 10; }
+};
+
+TEST_F(SchedulerTest, LinkedAllocaDependencies) {
+  default_selector Selector{};
+  if (Selector.select_device().is_host()) {
+    std::cerr << "Not run due to host-only environment\n";
+    return;
+  }
+
+  // 1. create two commands: alloca + alloca and link them
+  // 2. call Scheduler::GraphBuilder::getOrCreateAllocaForReq
+  detail::Requirement Req = getMockRequirement();
+
+  cl::sycl::queue Queue1;
+  cl::sycl::detail::QueueImplPtr Q1 = cl::sycl::detail::getSyclObjImpl(Queue1);
+
+  sycl::device HostDevice;
+  std::shared_ptr<detail::queue_impl> DefaultHostQueue(new detail::queue_impl(
+      detail::getSyclObjImpl(HostDevice), /*AsyncHandler=*/{},
+      /*PropList=*/{}));
+
+  std::shared_ptr<cl::sycl::detail::MemObjRecord> Record{
+      new cl::sycl::detail::MemObjRecord(DefaultHostQueue->getContextImplPtr(),
+                                         10)};
+
+  MemObjMock MemObj(Record);
+  Req.MSYCLMemObj = &MemObj;
+
+  cl::sycl::detail::AllocaCommand AllocaCmd1(DefaultHostQueue, Req, false);
+  Record->MAllocaCommands.push_back(&AllocaCmd1);
+
+  MockCommand DepCmd(DefaultHostQueue, Req);
+  MockCommand DepDepCmd(DefaultHostQueue, Req);
+  DepCmd.MDeps.push_back({&DepDepCmd, DepDepCmd.getRequirement(), &AllocaCmd1});
+  DepDepCmd.MUsers.insert(&DepCmd);
+  Record->MWriteLeaves.push_back(&DepCmd);
+
+  MockScheduler MS;
+  cl::sycl::detail::Command *AllocaCmd2 =
+      MS.getOrCreateAllocaForReq(Record.get(), &Req, Q1);
+
+  ASSERT_TRUE(!!AllocaCmd1.MLinkedAllocaCmd)
+      << "No link appeared in existing command";
+  ASSERT_EQ(AllocaCmd1.MLinkedAllocaCmd, AllocaCmd2) << "Invalid link appeared";
+  ASSERT_GT(AllocaCmd1.MUsers.count(AllocaCmd2), 0u)
+      << "New alloca isn't in users of the old one";
+  ASSERT_GT(AllocaCmd2->MDeps.size(), 1u)
+      << "No deps appeared in the new alloca";
+  ASSERT_GT(DepCmd.MUsers.count(AllocaCmd2), 0u)
+      << "No deps appeared for leaves of record (i.e. deps of existing alloca)";
+  ASSERT_TRUE(std::find_if(AllocaCmd2->MDeps.begin(), AllocaCmd2->MDeps.end(),
+                           [&](const cl::sycl::detail::DepDesc &Dep) -> bool {
+                             return Dep.MDepCommand == &AllocaCmd1;
+                           }) != AllocaCmd2->MDeps.end())
+      << "No deps for existing alloca appeared in new alloca";
+  ASSERT_TRUE(std::find_if(AllocaCmd2->MDeps.begin(), AllocaCmd2->MDeps.end(),
+                           [&](const cl::sycl::detail::DepDesc &Dep) -> bool {
+                             return Dep.MDepCommand == &DepCmd;
+                           }) != AllocaCmd2->MDeps.end())
+      << "No deps for leaves (deps of existing alloca) appeared in new alloca";
+}

sycl/unittests/scheduler/SchedulerTestUtils.hpp

@@ -91,6 +91,13 @@ class MockScheduler : public cl::sycl::detail::Scheduler {
                              cl::sycl::detail::BlockingT Blocking) {
     return GraphProcessor::enqueueCommand(Cmd, EnqueueResult, Blocking);
   }
+
+  cl::sycl::detail::AllocaCommandBase *
+  getOrCreateAllocaForReq(cl::sycl::detail::MemObjRecord *Record,
+                          const cl::sycl::detail::Requirement *Req,
+                          cl::sycl::detail::QueueImplPtr Queue) {
+    return MGraphBuilder.getOrCreateAllocaForReq(Record, Req, Queue);
+  }
 };
 
 void addEdge(cl::sycl::detail::Command *User, cl::sycl::detail::Command *Dep,