Make reduce a free function

ginkgo-project · Sep 13, 2021 · 6d581d8 · 6d581d8
1 parent c935c71
commit 6d581d8
Show file tree

Hide file tree

Showing 11 changed files with 81 additions and 50 deletions.
diff --git a/common/cuda_hip/components/reduction.hpp.inc b/common/cuda_hip/components/reduction.hpp.inc
@@ -237,12 +237,12 @@ __global__ __launch_bounds__(default_block_size) void reduce_add_array(
 template <typename ValueType>
 __global__
     __launch_bounds__(default_block_size) void reduce_add_array_with_existing_value(
-        size_type size, const ValueType *__restrict__ source,
-        ValueType *__restrict__ result)
+        size_type size, const ValueType* __restrict__ source,
+        ValueType* __restrict__ result)
 {
     __shared__ UninitializedArray<ValueType, default_block_size> block_sum;
-    reduce_array(size, source, static_cast<ValueType *>(block_sum),
-                 [](const ValueType &x, const ValueType &y) { return x + y; });
+    reduce_array(size, source, static_cast<ValueType*>(block_sum),
+                 [](const ValueType& x, const ValueType& y) { return x + y; });
 
     if (threadIdx.x == 0) {
         result[blockIdx.x] += block_sum[0];

diff --git a/core/base/array.cpp b/core/base/array.cpp
@@ -95,20 +95,23 @@ void Array<ValueType>::fill(const ValueType value)
 
 
 template <typename ValueType>
-void Array<ValueType>::reduce(ValueType *value) const
+void reduce(const Array<ValueType>& input_arr, ValueType* value)
 {
-    this->get_executor()->run(array::make_reduce_array(
-        this->get_const_data(), this->get_num_elems(), value));
+    auto exec = input_arr.get_executor();
+    exec->run(array::make_reduce_array(input_arr.get_const_data(),
+                                       input_arr.get_num_elems(), value));
 }
 
 
 template <typename ValueType>
-ValueType Array<ValueType>::reduce(ValueType init_value) const
+ValueType reduce(const Array<ValueType>& input_arr, const ValueType init_value)
 {
-    auto value = Array<ValueType>(this->get_executor(), {init_value});
-    this->get_executor()->run(array::make_reduce_array(
-        this->get_const_data(), this->get_num_elems(), value.get_data()));
-    return this->get_executor()->copy_val_to_host(value.get_data());
+    auto exec = input_arr.get_executor();
+    auto value = Array<ValueType>(exec, {init_value});
+    exec->run(array::make_reduce_array(input_arr.get_const_data(),
+                                       input_arr.get_num_elems(),
+                                       value.get_data()));
+    return exec->copy_val_to_host(value.get_data());
 }
 
 
@@ -118,13 +121,13 @@ GKO_INSTANTIATE_FOR_EACH_TEMPLATE_TYPE(GKO_DECLARE_ARRAY_FILL);
 
 
 #define GKO_DECLARE_ARRAY_REDUCE(_type) \
-    void Array<_type>::reduce(_type *value) const
+    void reduce(const Array<_type>& arr, _type* value)
 
 GKO_INSTANTIATE_FOR_EACH_TEMPLATE_TYPE(GKO_DECLARE_ARRAY_REDUCE);
 
 
 #define GKO_DECLARE_ARRAY_REDUCE2(_type) \
-    _type Array<_type>::reduce(_type val) const
+    _type reduce(const Array<_type>& arr, const _type val)
 
 GKO_INSTANTIATE_FOR_EACH_TEMPLATE_TYPE(GKO_DECLARE_ARRAY_REDUCE2);
 

diff --git a/core/components/reduce_array.hpp b/core/components/reduce_array.hpp
@@ -47,8 +47,8 @@ namespace kernels {
 
 #define GKO_DECLARE_REDUCE_ARRAY_KERNEL(ValueType)                  \
     void reduce_array(std::shared_ptr<const DefaultExecutor> exec,  \
-                      const ValueType *data, size_type num_entries, \
-                      ValueType *val)
+                      const ValueType* data, size_type num_entries, \
+                      ValueType* val)
 
 
 #define GKO_DECLARE_ALL_AS_TEMPLATES \

diff --git a/cuda/components/reduce_array.cu b/cuda/components/reduce_array.cu
@@ -56,12 +56,12 @@ namespace components {
 constexpr int default_block_size = 512;
 
 
-#include "common/components/reduction.hpp.inc"
+#include "common/cuda_hip/components/reduction.hpp.inc"
 
 
 template <typename ValueType>
 void reduce_array(std::shared_ptr<const DefaultExecutor> exec,
-                  const ValueType *array, size_type size, ValueType *val)
+                  const ValueType* array, size_type size, ValueType* val)
 {
     auto block_results_val = array;
     size_type grid_dim = size;

diff --git a/cuda/test/base/array.cu b/cuda/test/base/array.cu
@@ -92,3 +92,27 @@ TYPED_TEST(Array, CanCopyBackTemporaryCloneOnDifferentExecutor)
 
     this->assert_equal_to_original_x(this->x);
 }
+
+
+TYPED_TEST(Array, CanBeReduced)
+{
+    using T = TypeParam;
+    auto cuda = gko::CudaExecutor::create(0, this->exec);
+    auto arr = gko::Array<TypeParam>(cuda, I<T>{T(4), T(6)});
+    auto out = gko::Array<TypeParam>(cuda, I<T>{T(2)});
+    gko::reduce(arr, out.get_data());
+
+    out.set_executor(cuda->get_master());
+    ASSERT_EQ(out.get_data()[0], TypeParam{12});
+}
+
+
+TYPED_TEST(Array, CanBeReduced2)
+{
+    using T = TypeParam;
+    auto cuda = gko::CudaExecutor::create(0, this->exec);
+    auto arr = gko::Array<TypeParam>(cuda, I<T>{T(4), T(6)});
+    auto out = gko::reduce(arr, TypeParam{T(3)});
+
+    ASSERT_EQ(out, TypeParam{13});
+}
diff --git a/dpcpp/components/reduce_array.dp.cpp b/dpcpp/components/reduce_array.dp.cpp
@@ -44,15 +44,8 @@ namespace components {
 
 template <typename ValueType>
 void reduce_array(std::shared_ptr<const DefaultExecutor> exec,
-                  const ValueType *array, size_type n, ValueType *val)
-{
-    // exec->get_queue()->submit([&](sycl::handler &cgh) {
-    //     cgh.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> idx_id) {
-    //         const auto idx = idx_id[0];
-    //         array[idx] = val;
-    //     });
-    // });
-}
+                  const ValueType* array, size_type n,
+                  ValueType* val) GKO_NOT_IMPLEMENTED;
 
 GKO_INSTANTIATE_FOR_EACH_TEMPLATE_TYPE(GKO_DECLARE_REDUCE_ARRAY_KERNEL);
 

diff --git a/hip/components/reduce_array.hip.cpp b/hip/components/reduce_array.hip.cpp
@@ -58,12 +58,12 @@ namespace components {
 constexpr int default_block_size = 512;
 
 
-#include "common/components/reduction.hpp.inc"
+#include "common/cuda_hip/components/reduction.hpp.inc"
 
 
 template <typename ValueType>
 void reduce_array(std::shared_ptr<const DefaultExecutor> exec,
-                  const ValueType *array, size_type size, ValueType *val)
+                  const ValueType* array, size_type size, ValueType* val)
 {
     auto block_results_val = array;
     size_type grid_dim = size;

diff --git a/include/ginkgo/core/base/array.hpp b/include/ginkgo/core/base/array.hpp
@@ -487,20 +487,6 @@ class Array {
      */
     void fill(const value_type value);
 
-    /**
-     * Reduce the values in the array
-     *
-     * @return the reduced value
-     */
-    value_type reduce(value_type init_val = 0) const;
-
-    /**
-     * Reduce the values in the array
-     *
-     * @param the reduced value
-     */
-    void reduce(value_type *value) const;
-
     /**
      * Returns the number of elements in the Array.
      *
@@ -585,6 +571,31 @@ class Array {
 };
 
 
+/**
+ * Reduce the values in the array
+ *
+ * @tparam The type of the input data
+ *
+ * @param [in] input_arr the input array to be reduced
+ * @param [in] init_val the initial value
+ * @return the reduced value
+ */
+template <typename ValueType>
+ValueType reduce(const Array<ValueType>& input_arr,
+                 const ValueType init_val = 0);
+
+/**
+ * Reduce the values in the array
+ *
+ * @tparam The type of the input data
+ *
+ * @param [in] input_arr the input array to be reduced
+ * @param [out,in] value the reduced value
+ */
+template <typename ValueType>
+void reduce(const Array<ValueType>& input_arr, ValueType* value);
+
+
 namespace detail {
 
 

diff --git a/omp/components/reduce_array.cpp b/omp/components/reduce_array.cpp
@@ -41,7 +41,7 @@ namespace components {
 
 template <typename ValueType>
 void reduce_array(std::shared_ptr<const DefaultExecutor> exec,
-                  const ValueType *array, size_type n, ValueType *val)
+                  const ValueType* array, size_type n, ValueType* val)
 {
     ValueType out = *val;
 #pragma omp declare reduction(add:ValueType : omp_out = omp_out + omp_in)

diff --git a/reference/components/reduce_array.cpp b/reference/components/reduce_array.cpp
@@ -44,7 +44,7 @@ namespace components {
 
 template <typename ValueType>
 void reduce_array(std::shared_ptr<const DefaultExecutor> exec,
-                  const ValueType *array, size_type n, ValueType *val)
+                  const ValueType* array, size_type n, ValueType* val)
 {
     *val = std::accumulate(array, array + n, ValueType(*val));
 }

diff --git a/reference/test/base/array.cpp b/reference/test/base/array.cpp
@@ -78,18 +78,18 @@ TYPED_TEST(Array, CanBeFilledWithValue)
 
 TYPED_TEST(Array, CanBeReduced)
 {
-    TypeParam out = 0.0;
-    this->x.reduce(&out);
+    TypeParam out = 0.5;
+    gko::reduce(this->x, &out);
 
-    ASSERT_EQ(out, TypeParam{7});
+    ASSERT_EQ(out, TypeParam{7.5});
 }
 
 
 TYPED_TEST(Array, CanBeReduced2)
 {
-    auto out = this->x.reduce();
+    auto out = gko::reduce(this->x, TypeParam{2});
 
-    ASSERT_EQ(out, TypeParam{7});
+    ASSERT_EQ(out, TypeParam{9});
 }