[SYCL] Test for fix of linked alloca's deps

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/test/scheduler/LinkedAlloca.cpp

@@ -0,0 +1,85 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -I %sycl_source_dir %s -o %t.out
+// RUN: env SYCL_PRINT_EXECUTION_GRAPH=always %t.out
+// RUN: cat graph_7after_addHostAccessor.dot
+// RUN: cat graph_7after_addHostAccessor.dot | FileCheck %s
+
+#include <CL/sycl.hpp>
+
+namespace S = cl::sycl;
+
+void test() {
+  auto EH = [](S::exception_list EL) {
+    for (const std::exception_ptr &E : EL) {
+      throw E;
+    }
+  };
+
+  S::queue Queue(EH);
+
+  static const size_t BSIZE = 10;
+  S::buffer<int, 1> Buf{BSIZE};
+
+  // 0. submit kernel
+  Queue.submit([&](S::handler &CGH) {
+    S::accessor<int, 1, S::access::mode::write,
+                S::access::target::global_buffer>
+        GeneratorAcc(Buf, CGH);
+
+    auto GeneratorKernel = [GeneratorAcc]() {
+      for (size_t Idx = 0; Idx < GeneratorAcc.get_count(); ++Idx)
+        GeneratorAcc[Idx] = BSIZE - Idx;
+    };
+
+    CGH.single_task<class GeneratorTask0>(GeneratorKernel);
+  });
+  // 1. create host-accessor
+  {
+    S::accessor<int, 1, S::access::mode::write,
+                S::access::target::host_buffer>
+        Acc(Buf);
+
+    for (size_t Idx = 0; Idx < Acc.get_count(); ++Idx)
+      Acc[Idx] = -1;
+  }
+
+  // 2. submit kernel
+  Queue.submit([&](S::handler &CGH) {
+    S::accessor<int, 1, S::access::mode::write,
+                S::access::target::global_buffer>
+        GeneratorAcc(Buf, CGH);
+
+    auto GeneratorKernel = [GeneratorAcc]() {
+      for (size_t Idx = 0; Idx < GeneratorAcc.get_count(); ++Idx)
+        GeneratorAcc[Idx] = Idx;
+    };
+
+    CGH.single_task<class GeneratorTask>(GeneratorKernel);
+  });
+
+  // 3. create host-accessor
+  {
+    S::accessor<int, 1, S::access::mode::write,
+                S::access::target::host_buffer>
+        Acc(Buf);
+
+    bool Failure = false;
+    for (size_t Idx = 0; Idx < Acc.get_count(); ++Idx) {
+      fprintf(stderr, "Buffer [%03zu] = %i\n", Idx, Acc[Idx]);
+      Failure |= (Acc[Idx] != Idx);
+    }
+
+    assert(!Failure || "Invalid data in buffer");
+  }
+}
+
+int main(void) {
+  test();
+
+  return 0;
+}
+
+// CHECK: {{^}}"[[MAP_OP:0x[0-9a-fA-F]+]]"{{.*}} MAP ON
+// CHECK-DAG: "[[MAP_OP]]" -> "[[MAP_DEP_1:0x[0-9a-fA-F]+]]"
+// CHECK-DAG: "[[MAP_OP]]" -> "[[MAP_DEP_2:0x[0-9a-fA-F]+]]"
+// CHECK-DAG: {{^}}"[[MAP_DEP_1]]"{{.*}}\nEXEC CG ON
+// CHECK-DAG: {{^}}"[[MAP_DEP_2]]"{{.*}}\nALLOCA ON HOST

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=*/{},
+      detail::QueueOrder::Ordered, /*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,