Updates from Adam's comments (#1766)

* Updates from Adam's comments * Update for clang-format * Slight fix * Fix for `inclusive_scan` / `exclusive_scan` of size 0 * enabled back empty test in xclusive scans --------- Co-authored-by: Łukasz Ślusarczyk <lukasz.slusarczyk@intel.com>
uxlfoundation · Aug 7, 2024 · 68f59fe · 68f59fe
1 parent 3557043
commit 68f59fe
Show file tree

Hide file tree

Showing 3 changed files with 16 additions and 8 deletions.
diff --git a/include/oneapi/dpl/internal/distributed_ranges_impl/sp/algorithms/exclusive_scan.hpp b/include/oneapi/dpl/internal/distributed_ranges_impl/sp/algorithms/exclusive_scan.hpp
@@ -40,6 +40,12 @@ template <typename ExecutionPolicy, distributed_contiguous_range R, distributed_
 void
 exclusive_scan_impl_(ExecutionPolicy&& policy, R&& r, O&& o, U init, BinaryOp binary_op)
 {
+    if (stdrng::begin(r) == stdrng::end(r))
+    {
+        *stdrng::begin(o) = init;
+        return;
+    }
+
     using T = stdrng::range_value_t<O>;
 
     static_assert(std::is_same_v<std::remove_cvref_t<ExecutionPolicy>, sycl_device_collection>);
@@ -52,6 +58,7 @@ exclusive_scan_impl_(ExecutionPolicy&& policy, R&& r, O&& o, U init, BinaryOp bi
     std::vector<sycl::event> events;
 
     std::size_t segment_id = 0;
+
     for (auto&& segs : zipped_segments)
     {
         auto&& [in_segment, out_segment] = segs;
@@ -71,7 +78,7 @@ exclusive_scan_impl_(ExecutionPolicy&& policy, R&& r, O&& o, U init, BinaryOp bi
 
     std::vector<U> inits(stdrng::size(zipped_segments));
 
-    copy(d_inits, d_inits + inits.size(), inits.data() + 1);
+    copy(d_inits, d_inits + (inits.size() - 1), inits.data() + 1);
 
     sycl::free(d_inits, context());
 

diff --git a/include/oneapi/dpl/internal/distributed_ranges_impl/sp/algorithms/inclusive_scan.hpp b/include/oneapi/dpl/internal/distributed_ranges_impl/sp/algorithms/inclusive_scan.hpp
@@ -44,6 +44,11 @@ inclusive_scan_impl_(ExecutionPolicy&& policy, R&& r, O&& o, BinaryOp binary_op,
 {
     using T = stdrng::range_value_t<O>;
 
+    if (stdrng::begin(r) == stdrng::end(r))
+    {
+        return;
+    }
+
     static_assert(std::is_same_v<std::remove_cvref_t<ExecutionPolicy>, sycl_device_collection>);
 
     auto zipped_view = views::zip(r, o);
@@ -100,10 +105,7 @@ inclusive_scan_impl_(ExecutionPolicy&& policy, R&& r, O&& o, BinaryOp binary_op,
 
         auto e = q.submit([&](auto&& h) {
             h.depends_on(event);
-            h.single_task([=]() {
-                stdrng::range_value_t<O> value = *src_iter;
-                *dst_iter = value;
-            });
+            h.single_task([=]() { *dst_iter = *src_iter; });
         });
 
         events.push_back(e);
@@ -125,10 +127,9 @@ inclusive_scan_impl_(ExecutionPolicy&& policy, R&& r, O&& o, BinaryOp binary_op,
     std::size_t idx = 0;
     for (auto&& segs : zipped_segments)
     {
-        auto&& [in_segment, out_segment] = segs;
-
         if (idx > 0)
         {
+            auto&& [in_segment, out_segment] = segs;
             auto&& q = __detail::queue(ranges::rank(out_segment));
 
             auto first = stdrng::begin(out_segment);

diff --git a/test/distributed_ranges/sp/CMakeLists.txt b/test/distributed_ranges/sp/CMakeLists.txt
@@ -101,7 +101,7 @@ endfunction()
 # sort: https://github.com/oneapi-src/distributed-ranges/issues/785
 # ExclusiveScan.empty: https://github.com/oneapi-src/distributed-ranges/issues/791
 # IotaView.ZipWithDR: https://github.com/oneapi-src/distributed-ranges/issues/792
-set(test-exclusions Sort*:SpTests.Sort*:ExclusiveScan*empty:IotaView*ZipWithDR)
+set(test-exclusions Sort*:SpTests.Sort*:IotaView*ZipWithDR)
 
 add_sp_ctest(sp_tests sp_tests --gtest_filter=-${test-exclusions})
 add_sp_ctest(sp_tests_3 sp_tests --devicesCount 3 --gtest_filter=-${test-exclusions})