Use shp::devices() inside algorithms, disable non-aligned (#134)

* Use `shp::devices()` inside algorithms, disable non-aligned
`inclusive_scan` tests

* Use `shp::devices()` inside `inclusive_scan`

include/dr/shp/algorithms/for_each.hpp

@@ -25,11 +25,10 @@ void for_each(ExecutionPolicy &&policy, R &&r, Fn &&fn) {
   static_assert( // currently only one policy supported
       std::is_same_v<std::remove_cvref_t<ExecutionPolicy>, device_policy>);
 
-  auto &&devices = policy.get_devices();
   std::vector<sycl::event> events;
 
   for (auto &&segment : lib::ranges::segments(r)) {
-    auto device = devices[lib::ranges::rank(segment)];
+    auto device = shp::devices()[lib::ranges::rank(segment)];
 
     sycl::queue q(shp::context(), device);
 

include/dr/shp/algorithms/inclusive_scan.hpp

@@ -36,11 +36,9 @@ void inclusive_scan_impl_(ExecutionPolicy &&policy, R &&r, O &&o,
   if constexpr (std::is_same_v<std::remove_cvref_t<ExecutionPolicy>,
                                device_policy>) {
 
-    auto &&devices = std::forward<ExecutionPolicy>(policy).get_devices();
-
     std::vector<sycl::event> events;
 
-    auto root = devices[0];
+    auto root = shp::devices()[0];
     shp::device_allocator<T> allocator(shp::context(), root);
     shp::vector<T, shp::device_allocator<T>> partial_sums(
         std::size_t(zipped_segments.size()), allocator);
@@ -49,7 +47,7 @@ void inclusive_scan_impl_(ExecutionPolicy &&policy, R &&r, O &&o,
     for (auto &&segs : zipped_segments) {
       auto &&[in_segment, out_segment] = segs;
 
-      auto device = devices[lib::ranges::rank(in_segment)];
+      auto device = shp::devices()[lib::ranges::rank(in_segment)];
 
       sycl::queue q(shp::context(), device);
       oneapi::dpl::execution::device_policy local_policy(q);
@@ -115,7 +113,7 @@ void inclusive_scan_impl_(ExecutionPolicy &&policy, R &&r, O &&o,
     std::size_t idx = 0;
     for (auto &&segs : zipped_segments) {
       auto &&[in_segment, out_segment] = segs;
-      auto device = devices[lib::ranges::rank(out_segment)];
+      auto device = shp::devices()[lib::ranges::rank(out_segment)];
 
       sycl::queue q(shp::context(), device);
       oneapi::dpl::execution::device_policy local_policy(q);

include/dr/shp/algorithms/reduce.hpp

@@ -51,12 +51,10 @@ T reduce(ExecutionPolicy &&policy, R &&r, T init, BinaryOp &&binary_op) {
         lib::ranges::segments(r)[0].begin(), lib::ranges::segments(r)[0].end(),
         init, std::forward<BinaryOp>(binary_op)));
 
-    auto &&devices = std::forward<ExecutionPolicy>(policy).get_devices();
-
     std::vector<future_t> futures;
 
     for (auto &&segment : lib::ranges::segments(r)) {
-      auto device = devices[lib::ranges::rank(segment)];
+      auto device = shp::devices()[lib::ranges::rank(segment)];
 
       sycl::queue q(shp::context(), device);
       oneapi::dpl::execution::device_policy local_policy(q);

test/gtest/shp/algorithms.cpp

@@ -19,11 +19,10 @@ TEST(ShpTests, Iota) {
 }
 
 // hard to reproduce fails
-TEST(ShpTests, DISABLED_InclusiveScan) {
+TEST(ShpTests, InclusiveScan_aligned) {
   std::size_t n = 100;
 
   shp::distributed_vector<int, shp::device_allocator<int>> v(n);
-  shp::distributed_vector<int, shp::device_allocator<int>> o(v.size() * 2);
   std::vector<int> lv(n);
 
   // Range case, no binary op or init, perfectly aligned
@@ -39,44 +38,52 @@ TEST(ShpTests, DISABLED_InclusiveScan) {
     EXPECT_EQ(lv[i], v[i]);
   }
 
-  // Range case, binary op no init, non-aligned ranges
+  // Iterator case, no binary op or init, perfectly aligned
   for (auto &&x : lv) {
     x = lrand48() % 100;
   }
   shp::copy(lv.begin(), lv.end(), v.begin());
 
   std::inclusive_scan(lv.begin(), lv.end(), lv.begin());
-  shp::inclusive_scan(shp::par_unseq, v, o, std::plus<>());
+  shp::inclusive_scan(shp::par_unseq, v.begin(), v.end(), v.begin());
 
   for (std::size_t i = 0; i < lv.size(); i++) {
-    EXPECT_EQ(lv[i], o[i]);
+    EXPECT_EQ(lv[i], v[i]);
   }
+}
 
-  // Range case, binary op, init, non-aligned ranges
+TEST(ShpTests, DISABLED_InclusiveScan_nonaligned) {
+  std::size_t n = 100;
+
+  shp::distributed_vector<int, shp::device_allocator<int>> v(n);
+  shp::distributed_vector<int, shp::device_allocator<int>> o(v.size() * 2);
+  std::vector<int> lv(n);
+
+  // Range case, binary op no init, non-aligned ranges
   for (auto &&x : lv) {
     x = lrand48() % 100;
   }
   shp::copy(lv.begin(), lv.end(), v.begin());
 
-  std::inclusive_scan(lv.begin(), lv.end(), lv.begin(), std::multiplies<>(),
-                      12);
-  shp::inclusive_scan(shp::par_unseq, v, o, std::multiplies<>(), 12);
+  std::inclusive_scan(lv.begin(), lv.end(), lv.begin());
+  shp::inclusive_scan(shp::par_unseq, v, o, std::plus<>());
 
   for (std::size_t i = 0; i < lv.size(); i++) {
     EXPECT_EQ(lv[i], o[i]);
   }
 
-  // Iterator case, no binary op or init, perfectly aligned
+  // Range case, binary op, init, non-aligned ranges
   for (auto &&x : lv) {
     x = lrand48() % 100;
   }
   shp::copy(lv.begin(), lv.end(), v.begin());
 
-  std::inclusive_scan(lv.begin(), lv.end(), lv.begin());
-  shp::inclusive_scan(shp::par_unseq, v.begin(), v.end(), v.begin());
+  std::inclusive_scan(lv.begin(), lv.end(), lv.begin(), std::multiplies<>(),
+                      12);
+  shp::inclusive_scan(shp::par_unseq, v, o, std::multiplies<>(), 12);
 
   for (std::size_t i = 0; i < lv.size(); i++) {
-    EXPECT_EQ(lv[i], v[i]);
+    EXPECT_EQ(lv[i], o[i]);
   }
 
   // Iterator case, binary op no init, non-aligned ranges