mhp::exclusive_scan (#545)

benchmarks/gbench/common/inclusive_exclusive_scan.cpp

@@ -0,0 +1,100 @@
+// SPDX-FileCopyrightText: Intel Corporation
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#include "../common/dr_bench.hpp"
+
+using T = float;
+
+void Inclusive_Exclusive_Scan_DR(benchmark::State &state, bool is_inclusive) {
+  xhp::distributed_vector<T> a(default_vector_size, 3);
+  xhp::distributed_vector<T> b(default_vector_size, 0);
+  Stats stats(state, sizeof(T) * a.size(), sizeof(T) * b.size());
+
+  for (auto _ : state) {
+    for (std::size_t i = 0; i < default_repetitions; i++) {
+      stats.rep();
+      if (is_inclusive) {
+        xhp::inclusive_scan(a, b);
+      } else {
+        xhp::exclusive_scan(a, b, 0);
+      }
+    }
+  }
+}
+
+void Inclusive_Scan_DR(benchmark::State &state) {
+  Inclusive_Exclusive_Scan_DR(state, true);
+}
+DR_BENCHMARK(Inclusive_Scan_DR);
+
+void Exclusive_Scan_DR(benchmark::State &state) {
+  Inclusive_Exclusive_Scan_DR(state, false);
+}
+DR_BENCHMARK(Exclusive_Scan_DR);
+
+#ifdef SYCL_LANGUAGE_VERSION
+void Inclusive_Exclusive_Scan_Reference(benchmark::State &state,
+                                        bool is_inclusive) {
+  auto q = get_queue();
+  auto policy = oneapi::dpl::execution::make_device_policy(q);
+  auto a = sycl::malloc_device<T>(default_vector_size, q);
+  auto b = sycl::malloc_device<T>(default_vector_size, q);
+  Stats stats(state, sizeof(T) * default_vector_size,
+              sizeof(T) * default_vector_size);
+
+  for (auto _ : state) {
+    for (std::size_t i = 0; i < default_repetitions; i++) {
+      stats.rep();
+      if (is_inclusive) {
+        std::inclusive_scan(policy, a, a + default_vector_size, b,
+                            std::plus<T>{});
+      } else {
+        std::exclusive_scan(policy, a, a + default_vector_size, b, 0,
+                            std::plus<T>{});
+      }
+    }
+  }
+  sycl::free(a, q);
+  sycl::free(b, q);
+}
+
+void Inclusive_Scan_Reference(benchmark::State &state) {
+  Inclusive_Exclusive_Scan_Reference(state, true);
+}
+DR_BENCHMARK(Inclusive_Scan_Reference);
+
+void Exclusive_Scan_Reference(benchmark::State &state) {
+  Inclusive_Exclusive_Scan_Reference(state, false);
+}
+DR_BENCHMARK(Exclusive_Scan_Reference);
+#endif
+
+void Inclusive_Exclusive_Scan_Std(benchmark::State &state, bool is_inclusive) {
+  std::vector<T> a(default_vector_size, 3);
+  std::vector<T> b(default_vector_size, 0);
+  Stats stats(state, sizeof(T) * a.size(), sizeof(T) * b.size());
+
+  for (auto _ : state) {
+    for (std::size_t i = 0; i < default_repetitions; i++) {
+      stats.rep();
+      if (is_inclusive) {
+        std::inclusive_scan(std::execution::par_unseq, rng::begin(a),
+                            rng::end(a), rng::begin(b));
+      } else {
+        std::exclusive_scan(std::execution::par_unseq, rng::begin(a),
+                            rng::end(a), rng::begin(b), 0);
+      }
+    }
+  }
+}
+
+void Inclusive_Scan_Std(benchmark::State &state) {
+  Inclusive_Exclusive_Scan_Std(state, true);
+}
+DR_BENCHMARK(Inclusive_Scan_Std);
+
+void Exclusive_Scan_Std(benchmark::State &state) {
+  Inclusive_Exclusive_Scan_Std(state, false);
+}
+DR_BENCHMARK(Exclusive_Scan_Std);

benchmarks/gbench/mhp/CMakeLists.txt

@@ -10,7 +10,7 @@ add_executable(
   mhp-bench.cpp
   ../common/distributed_vector.cpp
   ../common/dot_product.cpp
-  ../common/inclusive_scan.cpp
+  ../common/inclusive_exclusive_scan.cpp
   ../common/sort.cpp
   ../common/stream.cpp
   wave_equation.cpp

benchmarks/gbench/shp/CMakeLists.txt

@@ -19,7 +19,7 @@ add_executable(
   gemm.cpp
   ../common/distributed_vector.cpp
   ../common/dot_product.cpp
-  ../common/inclusive_scan.cpp
+  ../common/inclusive_exclusive_scan.cpp
   ../common/sort.cpp
   ../common/stream.cpp)
 # cmake-format: on

include/dr/mhp.hpp

@@ -61,6 +61,7 @@
 #include <dr/mhp/algorithms/copy.hpp>
 #include <dr/mhp/algorithms/fill.hpp>
 #include <dr/mhp/algorithms/for_each.hpp>
+#include <dr/mhp/algorithms/exclusive_scan.hpp>
 #include <dr/mhp/algorithms/inclusive_scan.hpp>
 #include <dr/mhp/algorithms/iota.hpp>
 #include <dr/mhp/algorithms/reduce.hpp>

include/dr/mhp/algorithms/exclusive_scan.hpp

@@ -0,0 +1,49 @@
+// SPDX-FileCopyrightText: Intel Corporation
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#pragma once
+
+#include <dr/mhp/algorithms/inclusive_exclusive_scan_impl.hpp>
+
+namespace dr::mhp {
+
+template <dr::distributed_contiguous_range R,
+          dr::distributed_contiguous_range O, typename T, typename BinaryOp>
+auto exclusive_scan(R &&r, O &&o, T init, BinaryOp &&binary_op) {
+  return __detail::inclusive_exclusive_scan_impl_<true>(
+      std::forward<R>(r), rng::begin(std::forward<O>(o)),
+      std::forward<BinaryOp>(binary_op), std::optional(init));
+}
+
+template <dr::distributed_contiguous_range R,
+          dr::distributed_contiguous_range O, typename T>
+auto exclusive_scan(R &&r, O &&o, T init) {
+  return dr::mhp::exclusive_scan(std::forward<R>(r), std::forward<O>(o), init,
+                                 std::plus<rng::range_value_t<R>>());
+}
+
+// Distributed iterator versions
+
+template <dr::distributed_iterator Iter, dr::distributed_iterator OutputIter,
+          typename T, typename BinaryOp>
+OutputIter exclusive_scan(Iter first, Iter last, OutputIter d_first, T init,
+                          BinaryOp &&binary_op) {
+
+  return dr::mhp::exclusive_scan(rng::subrange(first, last), d_first,
+                                 std::forward<BinaryOp>(binary_op), init);
+}
+
+template <dr::distributed_iterator Iter, dr::distributed_iterator OutputIter,
+          typename T>
+OutputIter exclusive_scan(Iter first, Iter last, OutputIter d_first, T init) {
+  auto dist = rng::distance(first, last);
+  auto d_last = d_first;
+  rng::advance(d_last, dist);
+  dr::mhp::exclusive_scan(rng::subrange(first, last),
+                          rng::subrange(d_first, d_last), init);
+
+  return d_last;
+}
+
+} // namespace dr::mhp