Namespace/Docstring Fixes for Reduction (#10471)

This PR adds detail namespace for for simple/compound ops for reduction to make code more consistent with segmented reduction. It also includes minor docstring fixes.

cc @bdice

Authors:
  - Michael Wang (https://github.com/isVoid)

Approvers:
  - Mike Wilson (https://github.com/hyperbolic2346)
  - Bradley Dice (https://github.com/bdice)

URL: #10471

cpp/include/cudf/detail/reduction_functions.hpp

@@ -192,6 +192,8 @@ std::unique_ptr<scalar> mean(
  *
  * @param col input column to compute variance.
  * @param output_dtype data type of return type and typecast elements of input column.
+ * @param ddof Delta degrees of freedom. The divisor used is N - ddof, where N represents the number
+ * of elements.
  * @param stream CUDA stream used for device memory operations and kernel launches.
  * @param mr Device memory resource used to allocate the returned scalar's device memory.
  * @return Variance as scalar of type `output_dtype`.
@@ -213,6 +215,8 @@ std::unique_ptr<scalar> variance(
  *
  * @param col input column to compute standard deviation.
  * @param output_dtype data type of return type and typecast elements of input column.
+ * @param ddof Delta degrees of freedom. The divisor used is N - ddof, where N represents the number
+ * of elements.
  * @param stream CUDA stream used for device memory operations and kernel launches.
  * @param mr Device memory resource used to allocate the returned scalar's device memory.
  * @return Standard deviation as scalar of type `output_dtype`.

cpp/src/reductions/all.cu

@@ -88,11 +88,12 @@ std::unique_ptr<cudf::scalar> all(column_view const& col,
       dictionary_column_view(col).keys().type(), detail::all_fn{}, col, stream, mr);
   }
   // dispatch for non-dictionary types
-  return cudf::type_dispatcher(col.type(),
-                               simple::bool_result_element_dispatcher<cudf::reduction::op::min>{},
-                               col,
-                               stream,
-                               mr);
+  return cudf::type_dispatcher(
+    col.type(),
+    simple::detail::bool_result_element_dispatcher<cudf::reduction::op::min>{},
+    col,
+    stream,
+    mr);
 }
 
 }  // namespace reduction

cpp/src/reductions/any.cu

@@ -88,11 +88,12 @@ std::unique_ptr<cudf::scalar> any(column_view const& col,
       dictionary_column_view(col).keys().type(), detail::any_fn{}, col, stream, mr);
   }
   // dispatch for non-dictionary types
-  return cudf::type_dispatcher(col.type(),
-                               simple::bool_result_element_dispatcher<cudf::reduction::op::max>{},
-                               col,
-                               stream,
-                               mr);
+  return cudf::type_dispatcher(
+    col.type(),
+    simple::detail::bool_result_element_dispatcher<cudf::reduction::op::max>{},
+    col,
+    stream,
+    mr);
 }
 
 }  // namespace reduction

cpp/src/reductions/compound.cuh

@@ -25,22 +25,21 @@
 namespace cudf {
 namespace reduction {
 namespace compound {
+namespace detail {
 /**
- * @brief Multi-step reduction for operations such as mean and variance, and
- * standard deviation.
+ * @brief Multi-step reduction for operations such as mean, variance, and standard deviation.
  *
- * @param[in] col    input column view
- * @param[in] ddof   `Delta Degrees of Freedom` used for `std`, `var`.
- *                   The divisor used in calculations is N - ddof, where N
- *                   represents the number of elements.
- * @param[in] stream CUDA stream used for device memory operations and kernel launches.
- * @param[in] mr     Device memory resource used to allocate the returned scalar's device memory
- * @return    Output scalar in device memory
+ * @tparam ElementType  the input column data-type
+ * @tparam ResultType   the output data-type
+ * @tparam Op           the compound operator derived from `cudf::reduction::op::compound_op`
  *
- * @tparam ElementType  the input column cudf dtype
- * @tparam ResultType   the output cudf dtype
- * @tparam Op           the compound operator derived from
- * `cudf::reduction::op::compound_op`
+ * @param col input column view
+ * @param output_dtype data type of return type and typecast elements of input column.
+ * @param ddof Delta degrees of freedom used for standard deviation and variance. The divisor used
+ * is N - ddof, where N represents the number of elements.
+ * @param stream CUDA stream used for device memory operations and kernel launches.
+ * @param mr Device memory resource used to allocate the returned scalar's device memory
+ * @return Output scalar in device memory
  */
 template <typename ElementType, typename ResultType, typename Op>
 std::unique_ptr<scalar> compound_reduction(column_view const& col,
@@ -61,19 +60,19 @@ std::unique_ptr<scalar> compound_reduction(column_view const& col,
       auto it = thrust::make_transform_iterator(
         dcol->pair_begin<ElementType, true>(),
         compound_op.template get_null_replacing_element_transformer<ResultType>());
-      result = detail::reduce<Op, decltype(it), ResultType>(
+      result = cudf::reduction::detail::reduce<Op, decltype(it), ResultType>(
         it, col.size(), compound_op, valid_count, ddof, stream, mr);
     } else {
       auto it = thrust::make_transform_iterator(
         dcol->begin<ElementType>(), compound_op.template get_element_transformer<ResultType>());
-      result = detail::reduce<Op, decltype(it), ResultType>(
+      result = cudf::reduction::detail::reduce<Op, decltype(it), ResultType>(
         it, col.size(), compound_op, valid_count, ddof, stream, mr);
     }
   } else {
     auto it = thrust::make_transform_iterator(
       cudf::dictionary::detail::make_dictionary_pair_iterator<ElementType>(*dcol, col.has_nulls()),
       compound_op.template get_null_replacing_element_transformer<ResultType>());
-    result = detail::reduce<Op, decltype(it), ResultType>(
+    result = cudf::reduction::detail::reduce<Op, decltype(it), ResultType>(
       it, col.size(), compound_op, valid_count, ddof, stream, mr);
   }
 
@@ -152,6 +151,7 @@ struct element_type_dispatcher {
   }
 };
 
+}  // namespace detail
 }  // namespace compound
 }  // namespace reduction
 }  // namespace cudf

cpp/src/reductions/max.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -34,11 +34,12 @@ std::unique_ptr<cudf::scalar> max(column_view const& col,
   auto const dispatch_type = cudf::is_dictionary(col.type())
                                ? cudf::dictionary_column_view(col).indices().type()
                                : col.type();
-  return cudf::type_dispatcher(dispatch_type,
-                               simple::same_element_type_dispatcher<cudf::reduction::op::max>{},
-                               col,
-                               stream,
-                               mr);
+  return cudf::type_dispatcher(
+    dispatch_type,
+    simple::detail::same_element_type_dispatcher<cudf::reduction::op::max>{},
+    col,
+    stream,
+    mr);
 }
 
 }  // namespace reduction

cpp/src/reductions/mean.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -13,22 +13,27 @@
  * See the License for the specific language governing permissions and
  * limitations under the License.
  */
-// The translation unit for reduction `mean`
 
 #include <cudf/detail/reduction_functions.hpp>
 #include <cudf/dictionary/dictionary_column_view.hpp>
 #include <reductions/compound.cuh>
 
 #include <rmm/cuda_stream_view.hpp>
 
-std::unique_ptr<cudf::scalar> cudf::reduction::mean(column_view const& col,
-                                                    cudf::data_type const output_dtype,
-                                                    rmm::cuda_stream_view stream,
-                                                    rmm::mr::device_memory_resource* mr)
+namespace cudf {
+namespace reduction {
+
+std::unique_ptr<cudf::scalar> mean(column_view const& col,
+                                   cudf::data_type const output_dtype,
+                                   rmm::cuda_stream_view stream,
+                                   rmm::mr::device_memory_resource* mr)
 {
-  using reducer = cudf::reduction::compound::element_type_dispatcher<cudf::reduction::op::mean>;
+  using reducer = compound::detail::element_type_dispatcher<cudf::reduction::op::mean>;
   auto col_type =
     cudf::is_dictionary(col.type()) ? dictionary_column_view(col).keys().type() : col.type();
   return cudf::type_dispatcher(
     col_type, reducer(), col, output_dtype, /* ddof is not used for mean*/ 1, stream, mr);
 }
+
+}  // namespace reduction
+}  // namespace cudf

cpp/src/reductions/min.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -32,11 +32,12 @@ std::unique_ptr<cudf::scalar> min(column_view const& col,
   auto const dispatch_type = cudf::is_dictionary(col.type())
                                ? cudf::dictionary_column_view(col).indices().type()
                                : col.type();
-  return cudf::type_dispatcher(dispatch_type,
-                               simple::same_element_type_dispatcher<cudf::reduction::op::min>{},
-                               col,
-                               stream,
-                               mr);
+  return cudf::type_dispatcher(
+    dispatch_type,
+    simple::detail::same_element_type_dispatcher<cudf::reduction::op::min>{},
+    col,
+    stream,
+    mr);
 }
 
 }  // namespace reduction

cpp/src/reductions/product.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -30,7 +30,7 @@ std::unique_ptr<cudf::scalar> product(column_view const& col,
 {
   return cudf::type_dispatcher(
     cudf::is_dictionary(col.type()) ? dictionary_column_view(col).keys().type() : col.type(),
-    simple::element_type_dispatcher<cudf::reduction::op::product>{},
+    simple::detail::element_type_dispatcher<cudf::reduction::op::product>{},
     col,
     output_dtype,
     stream,

cpp/src/reductions/simple.cuh

@@ -38,6 +38,7 @@
 namespace cudf {
 namespace reduction {
 namespace simple {
+namespace detail {
 /**
  * @brief Reduction for 'sum', 'product', 'min', 'max', 'sum of squares'
  * which directly compute the reduction by a single step reduction call
@@ -64,11 +65,11 @@ std::unique_ptr<scalar> simple_reduction(column_view const& col,
     if (col.has_nulls()) {
       auto f  = simple_op.template get_null_replacing_element_transformer<ResultType>();
       auto it = thrust::make_transform_iterator(dcol->pair_begin<ElementType, true>(), f);
-      return detail::reduce(it, col.size(), simple_op, stream, mr);
+      return cudf::reduction::detail::reduce(it, col.size(), simple_op, stream, mr);
     } else {
       auto f  = simple_op.template get_element_transformer<ResultType>();
       auto it = thrust::make_transform_iterator(dcol->begin<ElementType>(), f);
-      return detail::reduce(it, col.size(), simple_op, stream, mr);
+      return cudf::reduction::detail::reduce(it, col.size(), simple_op, stream, mr);
     }
   }();
 
@@ -102,11 +103,11 @@ std::unique_ptr<scalar> fixed_point_reduction(column_view const& col,
     if (col.has_nulls()) {
       auto f  = simple_op.template get_null_replacing_element_transformer<Type>();
       auto it = thrust::make_transform_iterator(dcol->pair_begin<Type, true>(), f);
-      return detail::reduce(it, col.size(), simple_op, stream, mr);
+      return cudf::reduction::detail::reduce(it, col.size(), simple_op, stream, mr);
     } else {
       auto f  = simple_op.template get_element_transformer<Type>();
       auto it = thrust::make_transform_iterator(dcol->begin<Type>(), f);
-      return detail::reduce(it, col.size(), simple_op, stream, mr);
+      return cudf::reduction::detail::reduce(it, col.size(), simple_op, stream, mr);
     }
   }();
 
@@ -149,7 +150,7 @@ std::unique_ptr<scalar> dictionary_reduction(column_view const& col,
     auto p =
       cudf::dictionary::detail::make_dictionary_pair_iterator<ElementType>(*dcol, col.has_nulls());
     auto it = thrust::make_transform_iterator(p, f);
-    return detail::reduce(it, col.size(), simple_op, stream, mr);
+    return cudf::reduction::detail::reduce(it, col.size(), simple_op, stream, mr);
   }();
 
   // set scalar is valid
@@ -310,9 +311,9 @@ struct same_element_type_dispatcher {
                                      rmm::mr::device_memory_resource* mr)
   {
     if (!cudf::is_dictionary(col.type())) {
-      return simple::simple_reduction<ElementType, ElementType, Op>(col, stream, mr);
+      return simple_reduction<ElementType, ElementType, Op>(col, stream, mr);
     }
-    auto index = simple::simple_reduction<ElementType, ElementType, Op>(
+    auto index = simple_reduction<ElementType, ElementType, Op>(
       dictionary_column_view(col).get_indices_annotated(),
       stream,
       rmm::mr::get_current_device_resource());
@@ -442,6 +443,7 @@ struct element_type_dispatcher {
   }
 };
 
+}  // namespace detail
 }  // namespace simple
 }  // namespace reduction
 }  // namespace cudf

cpp/src/reductions/std.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -13,29 +13,26 @@
  * See the License for the specific language governing permissions and
  * limitations under the License.
  */
-// The translation unit for reduction `standard deviation`
 
 #include <cudf/detail/reduction_functions.hpp>
 #include <cudf/dictionary/dictionary_column_view.hpp>
 #include <reductions/compound.cuh>
 
 #include <rmm/cuda_stream_view.hpp>
 
-// @param[in] ddof Delta Degrees of Freedom used for `std`, `var`.
-//                 The divisor used in calculations is N - ddof, where N
-//                 represents the number of elements.
+namespace cudf {
+namespace reduction {
 
-std::unique_ptr<cudf::scalar> cudf::reduction::standard_deviation(
-  column_view const& col,
-  cudf::data_type const output_dtype,
-  cudf::size_type ddof,
-  rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
+std::unique_ptr<cudf::scalar> standard_deviation(column_view const& col,
+                                                 cudf::data_type const output_dtype,
+                                                 cudf::size_type ddof,
+                                                 rmm::cuda_stream_view stream,
+                                                 rmm::mr::device_memory_resource* mr)
 {
   // TODO: add cuda version check when the fix is available
 #if !defined(__CUDACC_DEBUG__)
   using reducer =
-    cudf::reduction::compound::element_type_dispatcher<cudf::reduction::op::standard_deviation>;
+    compound::detail::element_type_dispatcher<cudf::reduction::op::standard_deviation>;
   auto col_type =
     cudf::is_dictionary(col.type()) ? dictionary_column_view(col).keys().type() : col.type();
   return cudf::type_dispatcher(col_type, reducer(), col, output_dtype, ddof, stream, mr);
@@ -45,3 +42,6 @@ std::unique_ptr<cudf::scalar> cudf::reduction::standard_deviation(
   CUDF_FAIL("var/std reductions are not supported at debug build.");
 #endif
 }
+
+}  // namespace reduction
+}  // namespace cudf

cpp/src/reductions/sum.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -30,7 +30,7 @@ std::unique_ptr<cudf::scalar> sum(column_view const& col,
 {
   return cudf::type_dispatcher(
     cudf::is_dictionary(col.type()) ? dictionary_column_view(col).keys().type() : col.type(),
-    simple::element_type_dispatcher<cudf::reduction::op::sum>{},
+    simple::detail::element_type_dispatcher<cudf::reduction::op::sum>{},
     col,
     output_dtype,
     stream,

cpp/src/reductions/sum_of_squares.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -30,7 +30,7 @@ std::unique_ptr<cudf::scalar> sum_of_squares(column_view const& col,
 {
   return cudf::type_dispatcher(
     cudf::is_dictionary(col.type()) ? dictionary_column_view(col).keys().type() : col.type(),
-    simple::element_type_dispatcher<cudf::reduction::op::sum_of_squares>{},
+    simple::detail::element_type_dispatcher<cudf::reduction::op::sum_of_squares>{},
     col,
     output_dtype,
     stream,