[SYCL] Enqueue dependencies of blocked command

sycl/source/detail/scheduler/commands.hpp

@@ -493,10 +493,10 @@ class ExecCGCommand : public Command {
 
   detail::CG &getCG() const { return *MCommandGroup; }
 
-  // MEmptyCmd one is only employed if this command refers to host-task.
-  // MEmptyCmd due to unreliable mechanism of lookup for single EmptyCommand
-  // amongst users of host-task-representing command. This unreliability roots
-  // in cleanup process.
+  // MEmptyCmd is only employed if this command refers to host-task.
+  // The mechanism of lookup for single EmptyCommand amongst users of
+  // host-task-representing command is unreliable. This unreliability roots in
+  // the cleanup process.
   EmptyCommand *MEmptyCmd = nullptr;
 
 private:

sycl/source/detail/scheduler/graph_processor.cpp

@@ -58,19 +58,18 @@ bool Scheduler::GraphProcessor::enqueueCommand(Command *Cmd,
   if (!Cmd || Cmd->isSuccessfullyEnqueued())
     return true;
 
-  // Exit early if the command is blocked and the enqueue type is non-blocking
-  if (Cmd->isEnqueueBlocked() && !Blocking) {
-    EnqueueResult = EnqueueResultT(EnqueueResultT::SyclEnqueueBlocked, Cmd);
-    return false;
-  }
-
   // Recursively enqueue all the dependencies first and
   // exit immediately if any of the commands cannot be enqueued.
   for (DepDesc &Dep : Cmd->MDeps) {
     if (!enqueueCommand(Dep.MDepCommand, EnqueueResult, Blocking))
       return false;
   }
 
+  if (Cmd->isEnqueueBlocked() && !Blocking) {
+    EnqueueResult = EnqueueResultT(EnqueueResultT::SyclEnqueueBlocked, Cmd);
+    return false;
+  }
+
   return Cmd->enqueue(EnqueueResult, Blocking);
 }
 

sycl/test/host-interop-task/host-task-dependency2.cpp

@@ -0,0 +1,97 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// RUN: %CPU_RUN_PLACEHOLDER %t.out 10
+// RUN: %GPU_RUN_PLACEHOLDER %t.out 10
+// RUN: %ACC_RUN_PLACEHOLDER %t.out 10
+
+#include <CL/sycl.hpp>
+#include <iostream>
+
+using namespace cl::sycl;
+using namespace cl::sycl::access;
+
+static constexpr size_t BUFFER_SIZE = 1024;
+
+static auto EH = [](exception_list EL) {
+  for (const std::exception_ptr &E : EL) {
+    throw E;
+  }
+};
+
+// Host-task depending on another host-task via handler::depends_on() only
+// should not hang
+void test(size_t Count) {
+  queue Q(EH);
+
+  static constexpr size_t BufferSize = 10 * 1024;
+
+  buffer<int, 1> B0{range<1>{BufferSize}};
+  buffer<int, 1> B1{range<1>{BufferSize}};
+  buffer<int, 1> B2{range<1>{BufferSize}};
+  buffer<int, 1> B3{range<1>{BufferSize}};
+  buffer<int, 1> B4{range<1>{BufferSize}};
+  buffer<int, 1> B5{range<1>{BufferSize}};
+
+  for (size_t Idx = 1; Idx <= Count; ++Idx) {
+    // This host task should be submitted without hesitation
+    event E1 = Q.submit([&](handler &CGH) {
+      std::cout << "Submit 1" << std::endl;
+
+      auto Acc0 = B0.get_access<mode::read_write, target::host_buffer>(CGH);
+      auto Acc1 = B1.get_access<mode::read_write, target::host_buffer>(CGH);
+      auto Acc2 = B2.get_access<mode::read_write, target::host_buffer>(CGH);
+
+      CGH.codeplay_host_task([=] {
+        Acc0[0] = 1 * Idx;
+        Acc1[0] = 2 * Idx;
+        Acc2[0] = 3 * Idx;
+      });
+    });
+
+    // This host task is going to depend on blocked empty node of the first
+    // host-task (via buffer #2). Still this one should be enqueued.
+    event E2 = Q.submit([&](handler &CGH) {
+      std::cout << "Submit 2" << std::endl;
+
+      auto Acc2 = B2.get_access<mode::read_write, target::host_buffer>(CGH);
+      auto Acc3 = B3.get_access<mode::read_write, target::host_buffer>(CGH);
+
+      CGH.codeplay_host_task([=] {
+        Acc2[1] = 1 * Idx;
+        Acc3[1] = 2 * Idx;
+      });
+    });
+
+    // This host-task only depends on the second host-task via
+    // handler::depends_on(). This one should not hang and should be eexecuted
+    // after host-task #2.
+    event E3 = Q.submit([&](handler &CGH) {
+      CGH.depends_on(E2);
+
+      std::cout << "Submit 3" << std::endl;
+
+      auto Acc4 = B4.get_access<mode::read_write, target::host_buffer>(CGH);
+      auto Acc5 = B5.get_access<mode::read_write, target::host_buffer>(CGH);
+
+      CGH.codeplay_host_task([=] {
+        Acc4[2] = 1 * Idx;
+        Acc5[2] = 2 * Idx;
+      });
+    });
+  }
+
+  Q.wait_and_throw();
+}
+
+int main(int Argc, const char *Argv[]) {
+  size_t Count = 1;
+  if (Argc > 1)
+    Count = std::stoi(Argv[1]);
+
+  test(Count);
+  return 0;
+}

sycl/test/host-interop-task/host-task-dependency3.cpp

@@ -0,0 +1,126 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// RUN: %CPU_RUN_PLACEHOLDER %t.out 10
+// RUN: %GPU_RUN_PLACEHOLDER %t.out 10
+// RUN: %ACC_RUN_PLACEHOLDER %t.out 10
+
+#include <CL/sycl.hpp>
+#include <chrono>
+#include <iostream>
+#include <thread>
+
+using namespace cl::sycl;
+using namespace cl::sycl::access;
+
+static constexpr size_t BUFFER_SIZE = 1024;
+
+static auto EH = [](exception_list EL) {
+  for (const std::exception_ptr &E : EL) {
+    throw E;
+  }
+};
+
+// Host-task depending on another host-task via handler::depends_on() only
+// should not hang. A bit more complicated case with kernels depending on
+// host-task being involved.
+void test(size_t Count) {
+  queue Q(EH);
+
+  static constexpr size_t BufferSize = 10 * 1024;
+
+  buffer<int, 1> B0{range<1>{BufferSize}};
+  buffer<int, 1> B1{range<1>{BufferSize}};
+  buffer<int, 1> B2{range<1>{BufferSize}};
+  buffer<int, 1> B3{range<1>{BufferSize}};
+  buffer<int, 1> B4{range<1>{BufferSize}};
+  buffer<int, 1> B5{range<1>{BufferSize}};
+
+  using namespace std::chrono_literals;
+  constexpr auto SleepFor = 1s;
+
+  for (size_t Idx = 1; Idx <= Count; ++Idx) {
+    // This host task should be submitted without hesitation
+    Q.submit([&](handler &CGH) {
+      std::cout << "Submit HT-1" << std::endl;
+
+      auto Acc0 = B0.get_access<mode::read_write, target::host_buffer>(CGH);
+
+      CGH.codeplay_host_task([=] {
+        std::this_thread::sleep_for(SleepFor);
+        Acc0[0] = 1 * Idx;
+      });
+    });
+
+    Q.submit([&](handler &CGH) {
+      std::cout << "Submit Kernel-1" << std::endl;
+
+      auto Acc0 = B0.get_access<mode::read_write>(CGH);
+
+      CGH.single_task<class Test5_Kernel1>([=] { Acc0[1] = 1 * Idx; });
+    });
+
+    Q.submit([&](handler &CGH) {
+      std::cout << "Submit Kernel-2" << std::endl;
+
+      auto Acc1 = B1.get_access<mode::read_write>(CGH);
+
+      CGH.single_task<class Test5_Kernel2>([=] { Acc1[2] = 1 * Idx; });
+    });
+
+    Q.submit([&](handler &CGH) {
+      std::cout << "Submit HT-2" << std::endl;
+
+      auto Acc2 = B2.get_access<mode::read_write, target::host_buffer>(CGH);
+
+      CGH.codeplay_host_task([=] {
+        std::this_thread::sleep_for(SleepFor);
+        Acc2[3] = 1 * Idx;
+      });
+    });
+
+    // This host task is going to depend on blocked empty node of the second
+    // host-task (via buffer #0). Still this one should be enqueued.
+    event EHT3 = Q.submit([&](handler &CGH) {
+      std::cout << "Submit HT-3" << std::endl;
+
+      auto Acc0 = B0.get_access<mode::read_write, target::host_buffer>(CGH);
+      auto Acc1 = B1.get_access<mode::read_write, target::host_buffer>(CGH);
+      auto Acc2 = B2.get_access<mode::read_write, target::host_buffer>(CGH);
+
+      CGH.codeplay_host_task([=] {
+        std::this_thread::sleep_for(SleepFor);
+        Acc0[4] = 1 * Idx;
+        Acc1[4] = 2 * Idx;
+        Acc2[4] = 3 * Idx;
+      });
+    });
+
+    // This host-task only depends on the third host-task via
+    // handler::depends_on(). This one should not hang and should be executed
+    // after host-task #3.
+    Q.submit([&](handler &CGH) {
+      std::cout << "Submit HT-4" << std::endl;
+
+      CGH.depends_on(EHT3);
+
+      auto Acc5 = B5.get_access<mode::read_write, target::host_buffer>(CGH);
+
+      CGH.codeplay_host_task([=] { Acc5[5] = 1 * Idx; });
+    });
+  }
+
+  Q.wait_and_throw();
+}
+
+int main(int Argc, const char *Argv[]) {
+  size_t Count = 1;
+  if (Argc > 1)
+    Count = std::stoi(Argv[1]);
+
+  test(Count);
+  return 0;
+}

sycl/test/host-interop-task/host-task-dependency4.cpp

@@ -0,0 +1,30 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+#include <CL/sycl.hpp>
+
+cl::sycl::event submit(cl::sycl::queue &Q, cl::sycl::buffer<int> &B) {
+  return Q.submit([&](cl::sycl::handler &CGH) {
+    auto A = B.template get_access<cl::sycl::access::mode::read_write>(CGH);
+    CGH.codeplay_host_task([=]() { (void)A; });
+  });
+}
+
+int main() {
+  cl::sycl::queue Q;
+  int Status = 0;
+  cl::sycl::buffer<int> A{&Status, 1};
+  cl::sycl::vector_class<cl::sycl::event> Events;
+
+  Events.push_back(submit(Q, A));
+  Events.push_back(submit(Q, A));
+  Q.submit([&](sycl::handler &CGH) {
+     CGH.depends_on(Events);
+     CGH.codeplay_host_task([&] { printf("all done\n"); });
+   }).wait_and_throw();
+
+  return 0;
+}