[SYCL] Fix sync of host task vs kernel for in-order queue

sycl/source/detail/queue_impl.hpp

@@ -439,27 +439,41 @@ class queue_impl {
     return MAssertHappenedBuffer;
   }
 
-private:
-  void finalizeHandler(handler &Handler, const CG::CGTYPE &Type,
+protected:
+  // template is needed for proper unit testing
+  template <typename HandlerType = handler>
+  void finalizeHandler(HandlerType &Handler, const CG::CGTYPE &Type,
                        event &EventRet) {
     if (MIsInorder) {
-      bool NeedSeparateDependencyMgmt =
-          (Type == CG::CGTYPE::CodeplayHostTask ||
-           Type == CG::CGTYPE::CodeplayInteropTask);
+
+      auto IsExpDepManaged = [](const CG::CGTYPE &Type) {
+        return (Type == CG::CGTYPE::CodeplayHostTask ||
+                Type == CG::CGTYPE::CodeplayInteropTask);
+      };
+
       // Accessing and changing of an event isn't atomic operation.
       // Hence, here is the lock for thread-safety.
       std::lock_guard<std::mutex> Lock{MLastEventMtx};
 
+      if (MLastCGType == CG::CGTYPE::None)
+        MLastCGType = Type;
+      // Also handles case when sync model changes. E.g. Last is host, new is
+      // kernel.
+      bool NeedSeparateDependencyMgmt =
+          IsExpDepManaged(Type) || IsExpDepManaged(MLastCGType);
+
       if (NeedSeparateDependencyMgmt)
         Handler.depends_on(MLastEvent);
 
       EventRet = Handler.finalize();
 
       MLastEvent = EventRet;
+      MLastCGType = Type;
     } else
       EventRet = Handler.finalize();
   }
 
+private:
   /// Performs command group submission to the queue.
   ///
   /// \param CGF is a function object containing command group.
@@ -560,6 +574,10 @@ class queue_impl {
   // Access to the event should be guarded with MLastEventMtx
   event MLastEvent;
   std::mutex MLastEventMtx;
+  // Used for in-order queues in pair with MLastEvent
+  // Host tasks is explicitly synchronized in RT, pi tasks - implicitly by
+  // backend. Using type to setup explicit sync between host and pi tasks.
+  CG::CGTYPE MLastCGType = CG::CGTYPE::None;
 
   const bool MIsInorder;
 

sycl/unittests/scheduler/CMakeLists.txt

@@ -20,4 +20,5 @@ add_sycl_unittest(SchedulerTests OBJECT
     Regression.cpp
     utils.cpp
     LeafLimitDiffContexts.cpp
+    InOrderQueueSyncCheck.cpp
 )

sycl/unittests/scheduler/InOrderQueueSyncCheck.cpp

@@ -0,0 +1,106 @@
+//==---------- InOrderQueueSyncCheck.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"
+#include <CL/sycl.hpp>
+#include <detail/queue_impl.hpp>
+#include <detail/scheduler/commands.hpp>
+
+#include <gtest/gtest.h>
+
+using namespace sycl;
+
+// Define type with the only methods called by finalizeHandler
+class LimitedHandler {
+public:
+  virtual void depends_on(sycl::event){};
+
+  virtual event finalize() {
+    cl::sycl::detail::EventImplPtr NewEvent =
+        std::make_shared<detail::event_impl>();
+    return sycl::detail::createSyclObjFromImpl<sycl::event>(NewEvent);
+  };
+};
+
+// Needed to use EXPECT_CALL to verify depends_on that originally appends lst
+// event as dependency to the new CG
+class LimitedHandlerSimulation : public LimitedHandler {
+public:
+  MOCK_METHOD1(depends_on, void(sycl::event));
+};
+
+class MockQueueImpl : public sycl::detail::queue_impl {
+public:
+  MockQueueImpl(const sycl::detail::DeviceImplPtr &Device,
+                const sycl::async_handler &AsyncHandler,
+                const sycl::property_list &PropList)
+      : sycl::detail::queue_impl(Device, AsyncHandler, PropList) {}
+  using sycl::detail::queue_impl::finalizeHandler;
+};
+
+// Only check events dependency in queue_impl::finalizeHandler
+TEST_F(SchedulerTest, InOrderQueueSyncCheck) {
+  sycl::platform Plt{sycl::default_selector()};
+  if (Plt.is_host() || Plt.get_backend() == sycl::backend::ext_oneapi_cuda ||
+      Plt.get_backend() == sycl::backend::ext_oneapi_hip) {
+    std::cerr << "Test is not supported on "
+              << Plt.get_info<sycl::info::platform::name>() << ", skipping\n";
+    GTEST_SKIP(); // test is not supported on selected platform.
+  }
+
+  const sycl::device Dev = Plt.get_devices()[0];
+  auto Queue = std::make_shared<MockQueueImpl>(
+      sycl::detail::getSyclObjImpl(Dev), sycl::async_handler{},
+      sycl::property::queue::in_order());
+
+  // What we are testing here:
+  // Task type  | Must depend on
+  //  host      | yes - always, separate sync management
+  //  host      | yes - always, separate sync management
+  //  kernel    | yes - change of sync approach
+  //  kernel    | no  - sync between pi calls must be done by backend
+  //  host      | yes - always, separate sync management
+
+  sycl::event Event;
+  // host task
+  {
+    LimitedHandlerSimulation MockCGH;
+    EXPECT_CALL(MockCGH, depends_on).Times(1);
+    Queue->finalizeHandler<LimitedHandlerSimulation>(
+        MockCGH, detail::CG::CGTYPE::CodeplayHostTask, Event);
+  }
+  // host task
+  {
+    LimitedHandlerSimulation MockCGH;
+    EXPECT_CALL(MockCGH, depends_on).Times(1);
+    Queue->finalizeHandler<LimitedHandlerSimulation>(
+        MockCGH, detail::CG::CGTYPE::CodeplayHostTask, Event);
+  }
+  // kernel task
+  {
+    LimitedHandlerSimulation MockCGH;
+    EXPECT_CALL(MockCGH, depends_on).Times(1);
+    Queue->finalizeHandler<LimitedHandlerSimulation>(
+        MockCGH, detail::CG::CGTYPE::Kernel, Event);
+  }
+  // kernel task
+  {
+    LimitedHandlerSimulation MockCGH;
+    EXPECT_CALL(MockCGH, depends_on).Times(0);
+    Queue->finalizeHandler<LimitedHandlerSimulation>(
+        MockCGH, detail::CG::CGTYPE::Kernel, Event);
+  }
+  // host task
+  {
+    LimitedHandlerSimulation MockCGH;
+    EXPECT_CALL(MockCGH, depends_on).Times(1);
+    Queue->finalizeHandler<LimitedHandlerSimulation>(
+        MockCGH, detail::CG::CGTYPE::CodeplayHostTask, Event);
+  }
+}