[OpenCL] Handle synchronization between queues from different platforms

include/hipSYCL/runtime/cuda/cuda_queue.hpp

@@ -108,7 +108,7 @@ class cuda_queue : public inorder_queue
 
   /// Causes the queue to wait until an event on another queue has occured.
   /// the other queue must be from the same backend
-  virtual result submit_queue_wait_for(std::shared_ptr<dag_node_event> evt) override;
+  virtual result submit_queue_wait_for(dag_node_ptr evt) override;
   virtual result submit_external_wait_for(dag_node_ptr node) override;
 
   virtual result wait() override;

include/hipSYCL/runtime/hip/hip_queue.hpp

@@ -101,7 +101,7 @@ class hip_queue : public inorder_queue
 
   /// Causes the queue to wait until an event on another queue has occured.
   /// the other queue must be from the same backend
-  virtual result submit_queue_wait_for(std::shared_ptr<dag_node_event> evt) override;
+  virtual result submit_queue_wait_for(dag_node_ptr evt) override;
   virtual result submit_external_wait_for(dag_node_ptr node) override;
 
   virtual result wait() override;

include/hipSYCL/runtime/inorder_queue.hpp

@@ -69,7 +69,7 @@ class inorder_queue
 
   /// Causes the queue to wait until an event on another queue has occured.
   /// the other queue must be from the same backend
-  virtual result submit_queue_wait_for(std::shared_ptr<dag_node_event> evt) = 0;
+  virtual result submit_queue_wait_for(dag_node_ptr evt) = 0;
   virtual result submit_external_wait_for(dag_node_ptr node) = 0;
 
   virtual result wait() = 0;

include/hipSYCL/runtime/ocl/ocl_hardware_manager.hpp

@@ -103,6 +103,7 @@ class ocl_hardware_manager : public backend_hardware_manager
 
   cl::Platform get_platform(int platform_id);
   cl::Context get_context(int platform_id);
+  cl::Context get_context(device_id dev);
 private:
   std::vector<ocl_hardware_context> _devices;
   std::vector<cl::Platform> _platforms;

include/hipSYCL/runtime/ocl/ocl_queue.hpp

@@ -64,7 +64,7 @@ class ocl_queue : public inorder_queue
 
   /// Causes the queue to wait until an event on another queue has occured.
   /// the other queue must be from the same backend
-  virtual result submit_queue_wait_for(std::shared_ptr<dag_node_event> evt) override;
+  virtual result submit_queue_wait_for(dag_node_ptr evt) override;
   virtual result submit_external_wait_for(dag_node_ptr node) override;
 
   virtual result wait() override;

include/hipSYCL/runtime/omp/omp_queue.hpp

@@ -53,7 +53,7 @@ class omp_queue : public inorder_queue
 
   /// Causes the queue to wait until an event on another queue has occured.
   /// the other queue must be from the same backend
-  virtual result submit_queue_wait_for(std::shared_ptr<dag_node_event> evt) override;
+  virtual result submit_queue_wait_for(dag_node_ptr evt) override;
   virtual result submit_external_wait_for(dag_node_ptr node) override;
 
   virtual result wait() override;

include/hipSYCL/runtime/serialization/serialization.hpp

@@ -40,6 +40,7 @@ namespace hipsycl::rt {
 std::ostream &operator<<(std::ostream &out, const hardware_platform value);
 std::ostream &operator<<(std::ostream &out, const api_platform value);
 std::ostream &operator<<(std::ostream &out, const backend_id value);
+std::ostream &operator<<(std::ostream &out, device_id dev);
 
 template <int Dim>
 std::ostream &operator<<(std::ostream &out, const static_array<Dim> &v) {

include/hipSYCL/runtime/ze/ze_queue.hpp

@@ -35,6 +35,7 @@
 #include "../inorder_queue.hpp"
 #include "hipSYCL/runtime/code_object_invoker.hpp"
 #include "hipSYCL/runtime/event.hpp"
+#include "hipSYCL/runtime/hints.hpp"
 #include "ze_code_object.hpp"
 
 
@@ -62,7 +63,7 @@ class ze_queue : public inorder_queue
 
   /// Causes the queue to wait until an event on another queue has occured.
   /// the other queue must be from the same backend
-  virtual result submit_queue_wait_for(std::shared_ptr<dag_node_event> evt) override;
+  virtual result submit_queue_wait_for(dag_node_ptr evt) override;
   virtual result submit_external_wait_for(dag_node_ptr node) override;
 
   virtual result wait() override;

src/runtime/cuda/cuda_queue.cpp

@@ -411,7 +411,8 @@ result cuda_queue::submit_memset(memset_operation &op, dag_node_ptr node) {
 
 /// Causes the queue to wait until an event on another queue has occured.
 /// the other queue must be from the same backend
-result cuda_queue::submit_queue_wait_for(std::shared_ptr<dag_node_event> evt) {
+result cuda_queue::submit_queue_wait_for(dag_node_ptr node) {
+  auto evt = node->get_event();
   assert(dynamic_is<inorder_queue_event<cudaEvent_t>>(evt.get()));
 
   inorder_queue_event<cudaEvent_t> *cuda_evt =

src/runtime/hip/hip_queue.cpp

@@ -441,7 +441,8 @@ result hip_queue::query_status(inorder_queue_status &status) {
 
 /// Causes the queue to wait until an event on another queue has occured.
 /// the other queue must be from the same backend
-result hip_queue::submit_queue_wait_for(std::shared_ptr<dag_node_event> evt) {
+result hip_queue::submit_queue_wait_for(dag_node_ptr node) {
+  auto evt = node->get_event();
   assert(dynamic_is<inorder_queue_event<hipEvent_t>>(evt.get()));
 
   inorder_queue_event<hipEvent_t> *hip_evt =

src/runtime/inorder_executor.cpp

@@ -172,7 +172,7 @@ void inorder_executor::submit_directly(dag_node_ptr node, operation *op,
                    "requirement follows in the same inorder queue)"
                 << std::endl;
           } else {
-            res = _q->submit_queue_wait_for(req->get_event());
+            res = _q->submit_queue_wait_for(req);
           }
         }
       }

src/runtime/ocl/ocl_hardware_manager.cpp

@@ -421,7 +421,10 @@ ocl_hardware_manager::ocl_hardware_manager()
     int platform_id = _platforms.size() - 1;
 
     std::vector<cl::Device> devs;
-    err = p.getDevices(CL_DEVICE_TYPE_ALL, &devs);
+    // CL param validation layer does not like CL_DEVICE_TYPE_ALL here
+    err = p.getDevices(CL_DEVICE_TYPE_CPU | CL_DEVICE_TYPE_GPU |
+                           CL_DEVICE_TYPE_ACCELERATOR,
+                       &devs);
     if(err != CL_SUCCESS) {
       print_warning(
           __hipsycl_here(),
@@ -470,5 +473,10 @@ cl::Context ocl_hardware_manager::get_context(int platform_id) {
   return _platform_contexts[platform_id];
 }
 
+cl::Context ocl_hardware_manager::get_context(device_id dev) {
+  int platform_id = _devices[dev.get_id()].get_platform_id();
+  return _platform_contexts[platform_id];
+}
+
 }
 }

src/runtime/ocl/ocl_queue.cpp

@@ -27,6 +27,7 @@
 
 
 #include "hipSYCL/runtime/error.hpp"
+#include "hipSYCL/runtime/serialization/serialization.hpp"
 #include "hipSYCL/runtime/kernel_cache.hpp"
 #include "hipSYCL/runtime/inorder_queue.hpp"
 #include "hipSYCL/runtime/executor.hpp"
@@ -156,6 +157,13 @@ std::shared_ptr<dag_node_event> ocl_queue::create_queue_completion_event() {
 }
 
 result ocl_queue::submit_memcpy(memcpy_operation &op, dag_node_ptr) {
+
+  HIPSYCL_DEBUG_INFO << "ocl_queue: On device "
+                     << _hw_manager->get_device_id(_device_index)
+                     << ": Processing memcpy request from device "
+                     << op.source().get_device() << " to "
+                     << op.dest().get_device() << std::endl;
+
   // TODO We could probably unify some of the logic here between
   // backends
   device_id source_dev = op.source().get_device();
@@ -258,10 +266,18 @@ result ocl_queue::submit_memset(memset_operation& op, dag_node_ptr) {
 
 /// Causes the queue to wait until an event on another queue has occured.
 /// the other queue must be from the same backend
-result ocl_queue::submit_queue_wait_for(std::shared_ptr<dag_node_event> evt) {
-  ocl_node_event* ocl_evt = static_cast<ocl_node_event*>(evt.get());
+result ocl_queue::submit_queue_wait_for(dag_node_ptr evt) {
+
+  ocl_node_event *ocl_evt =
+      static_cast<ocl_node_event *>(evt->get_event().get());
+
   std::vector<cl::Event> events{ocl_evt->get_event()};
 
+  if (_hw_manager->get_context(ocl_evt->get_device()) !=
+      _hw_manager->get_context(_hw_manager->get_device_id(_device_index))) {
+    return submit_external_wait_for(evt);
+  }
+
   cl::Event wait_evt;
   cl_int err = _queue.enqueueBarrierWithWaitList(&events, &wait_evt);
 

src/runtime/omp/omp_queue.cpp

@@ -364,8 +364,9 @@ result omp_queue::submit_memset(memset_operation & op, dag_node_ptr node) {
 
   /// Causes the queue to wait until an event on another queue has occured.
   /// the other queue must be from the same backend
-result omp_queue::submit_queue_wait_for(std::shared_ptr<dag_node_event> evt) {
+result omp_queue::submit_queue_wait_for(dag_node_ptr node) {
   HIPSYCL_DEBUG_INFO << "omp_queue: Submitting wait for other queue..." << std::endl;
+  auto evt = node->get_event();
   if(!evt) {
     return register_error(
         __hipsycl_here(),

src/runtime/serialization/serialization.cpp

@@ -185,4 +185,9 @@ void device_id::dump(std::ostream &ostr) const {
   ostr << _backend.hw_platform << "-Device" << _device_id;
 }
 
+std::ostream &operator<<(std::ostream &out, device_id dev) {
+  dev.dump(out);
+  return out;
+}
+
 } // end of namespace hipsycl::rt

src/runtime/ze/ze_queue.cpp

@@ -37,6 +37,7 @@
 #include "hipSYCL/runtime/device_id.hpp"
 #include "hipSYCL/runtime/error.hpp"
 #include "hipSYCL/runtime/event.hpp"
+#include "hipSYCL/runtime/hints.hpp"
 #include "hipSYCL/runtime/inorder_queue.hpp"
 #include "hipSYCL/runtime/ze/ze_code_object.hpp"
 #include "hipSYCL/runtime/ze/ze_queue.hpp"
@@ -356,7 +357,8 @@ result ze_queue::wait() {
   return make_success();
 }
 
-result ze_queue::submit_queue_wait_for(std::shared_ptr<dag_node_event> evt) {
+result ze_queue::submit_queue_wait_for(dag_node_ptr node) {
+  auto evt = node->get_event();
   _enqueued_synchronization_ops.push_back(evt);
   return make_success();
 }