Fix device memory when using local in row distribution

Author: Mikolaj Komar

include/dr/mp/algorithms/reduce.hpp

@@ -35,11 +35,10 @@ inline auto dpl_reduce(rng::forward_range auto &&r, auto &&binary_op) {
                          sycl::known_identity_v<Fn, T>, binary_op);
     } else {
       dr::drlog.debug("  peel 1st value\n");
-      auto base = *rng::begin(r);
       return std::reduce(dpl_policy(),
                          dr::__detail::direct_iterator(rng::begin(r) + 1),
                          dr::__detail::direct_iterator(rng::end(r)),
-                         sycl_get(base), binary_op);
+                         sycl_get_deref(rng::begin(r)), binary_op);
     }
   }
 #else

include/dr/mp/containers/matrix_formats/csr_row_distribution.hpp

@@ -295,7 +295,12 @@ class csr_row_distribution {
     auto my_tuple = std::make_tuple(real_row_size, segment_size_ * rank, offset, local_rows);
     view_helper_const = alloc.allocate(1);
 
-    view_helper_const[0] = my_tuple;
+
+    if (use_sycl()) {
+      sycl_queue().memcpy(view_helper_const, &my_tuple, sizeof(view_tuple)).wait();
+    } else {
+      view_helper_const[0] = my_tuple;
+    }
 
     local_view = std::make_shared<view_type>(get_elem_view(vals_size_, view_helper_const, cols_data_, vals_data_, rank));
   }

include/dr/mp/sycl_support.hpp

@@ -17,6 +17,20 @@ sycl::queue &sycl_queue();
 
 namespace dr::mp::__detail {
 
+//sometimes we only want to dereference iterator inside SYCL
+template <typename T> auto sycl_get_deref(T v) {
+  using deref_type = decltype(*v);
+  deref_type temp;
+  {
+    sycl::buffer<deref_type> buff(&temp, 1);
+    sycl_queue().submit([&](auto &&h) {
+      sycl::accessor access(buff, h, sycl::write_only, sycl::no_init);
+      h.single_task([=](auto i) { access[0] = *v;});
+    }).wait();
+  }
+  return temp;
+}
+
 template <typename T> T sycl_get(T &v) {
   T temp;
   sycl_queue().memcpy(&temp, &v, sizeof(v)).wait();