prod, max, min, any, all scaffolding

cpp/CMakeLists.txt

@@ -367,7 +367,12 @@ add_library(
   src/reductions/scan/scan_inclusive.cu
   src/reductions/std.cu
   src/reductions/sum.cu
+  src/reductions/segmented_all.cu
+  src/reductions/segmented_any.cu
   src/reductions/segmented_sum.cu
+  src/reductions/segmented_product.cu
+  src/reductions/segmented_max.cu
+  src/reductions/segmented_min.cu
   src/reductions/sum_of_squares.cu
   src/reductions/var.cu
   src/replace/clamp.cu

cpp/include/cudf/detail/reduction.cuh

@@ -287,10 +287,15 @@ std::unique_ptr<column> segmented_reduce(InputIterator d_in,
 
 template <typename Op,
           typename InputIterator,
+          typename OffsetIterator,
           typename OutputType = typename thrust::iterator_value<InputIterator>::type,
-          typename std::enable_if_t<std::is_same_v<OutputType, string_view>>* = nullptr,
-          typename... Args>
-std::unique_ptr<column> segmented_reduce(Args&&...)
+          typename std::enable_if_t<std::is_same_v<OutputType, string_view>>* = nullptr>
+std::unique_ptr<column> segmented_reduce(InputIterator,
+                                         OffsetIterator,
+                                         cudf::size_type,
+                                         op::simple_op<Op>,
+                                         rmm::cuda_stream_view,
+                                         rmm::mr::device_memory_resource*)
 {
   CUDF_FAIL("Segment reduction for string type is unsupported.");
 }

cpp/include/cudf/detail/reduction_functions.hpp

@@ -255,19 +255,19 @@ std::unique_ptr<scalar> nth_element(
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Computes sums to each segment of the input column
+ * @brief Compute sum of each segment in input column
  *
- * If all elements in an input segment are null, the result is null.
+ * If all elements in input segment are null, the segment result is null.
  *
  * @throw cudf::logic_error if input column type is not convertible to `output_dtype`
  * @throw cudf::logic_error if `output_dtype` is not arithmetic point type
  *
- * @param col input column to compute sum
+ * @param col Input column to compute sum
  * @param offsets Indices to identify segment boundaries
- * @param output_dtype data type of return type and typecast elements of input column
+ * @param output_dtype Data type of return type and typecast elements of input column
  * @param stream CUDA stream used for device memory operations and kernel launches.
  * @param mr Device memory resource used to allocate the returned column's device memory
- * @return Sum as column of type `output_dtype`.
+ * @return Sums of segments in type `output_dtype`.
  */
 std::unique_ptr<column> segmented_sum(
   column_view const& col,
@@ -276,103 +276,112 @@ std::unique_ptr<column> segmented_sum(
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
-// /**
-//  * @brief Computes minimum of elements in input column
-//  *
-//  * If all elements in input column are null, output scalar is null.
-//  *
-//  * @throw cudf::logic_error if input column type is convertible to `output_dtype`
-//  *
-//  * @param col input column to compute minimum.
-//  * @param output_dtype data type of return type and typecast elements of input column
-//  * @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 Minimum element as scalar of type `output_dtype`.
-//  */
-// std::unique_ptr<scalar> min(
-//   column_view const& col,
-//   data_type const output_dtype,
-//   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
-//   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+/**
+ * @brief Computes product of elements in input column
+ *
+ * If all elements in input segment are null, the segment result is null.
+ *
+ * @throw cudf::logic_error if input column type is not convertible to `output_dtype`
+ * @throw cudf::logic_error if `output_dtype` is not arithmetic point type
+ *
+ * @param col input column to compute product.
+ * @param output_dtype data type of return type and typecast elements of input column
+ * @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 Product as scalar of type `output_dtype`.
+ */
+std::unique_ptr<column> segmented_product(
+  column_view const& col,
+  column_view const& offsets,
+  data_type const output_dtype,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
-// /**
-//  * @brief Computes maximum of elements in input column
-//  *
-//  * If all elements in input column are null, output scalar is null.
-//  *
-//  * @throw cudf::logic_error if input column type is convertible to `output_dtype`
-//  *
-//  * @param col input column to compute maximum.
-//  * @param output_dtype data type of return type and typecast elements of input column
-//  * @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 Maximum element as scalar of type `output_dtype`.
-//  */
-// std::unique_ptr<scalar> max(
-//   column_view const& col,
-//   data_type const output_dtype,
-//   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
-//   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+/**
+ * @brief Compute minimum of each segment in input column
+ *
+ * If all elements in input segment are null, the segment result is null.
+ *
+ * @throw cudf::logic_error if input column type is convertible to `output_dtype`
+ *
+ * @param col Input column to compute minimum.
+ * @param offsets Indices to identify segment boundaries
+ * @param output_dtype Data type of return type and typecast elements of input column
+ * @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 Minimums of segments in type `output_dtype`.
+ */
+std::unique_ptr<column> segmented_min(
+  column_view const& col,
+  column_view const& offsets,
+  data_type const output_dtype,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
-// /**
-//  * @brief Computes any of elements in input column is true when typecasted to bool
-//  *
-//  * If all elements in input column are null, output scalar is null.
-//  *
-//  * @throw cudf::logic_error if input column type is not convertible to bool
-//  * @throw cudf::logic_error if `output_dtype` is not bool
-//  *
-//  * @param col input column to compute any_of.
-//  * @param output_dtype data type of return type and typecast elements of input column
-//  * @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 bool scalar if any of elements is true when typecasted to bool
-//  */
-// std::unique_ptr<scalar> any(
-//   column_view const& col,
-//   data_type const output_dtype,
-//   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
-//   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+/**
+ * @brief Compute maximum of each segment in input column
+ *
+ * If all elements in input segment are null, the segment result is null.
+ *
+ * @throw cudf::logic_error if input column type is convertible to `output_dtype`
+ *
+ * @param col Input column to compute maximum.
+ * @param offsets Indices to identify segment boundaries
+ * @param output_dtype Data type of return type and typecast elements of input column
+ * @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 Maximums of segments in type `output_dtype`.
+ */
+std::unique_ptr<column> segmented_max(
+  column_view const& col,
+  column_view const& offsets,
+  data_type const output_dtype,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
-// /**
-//  * @brief Computes all of elements in input column is true when typecasted to bool
-//  *
-//  * If all elements in input column are null, output scalar is null.
-//  *
-//  * @throw cudf::logic_error if input column type is not convertible to bool
-//  * @throw cudf::logic_error if `output_dtype` is not bool
-//  *
-//  * @param col input column to compute all_of.
-//  * @param output_dtype data type of return type and typecast elements of input column
-//  * @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 bool scalar if all of elements is true when typecasted to bool
-//  */
-// std::unique_ptr<scalar> all(
-//   column_view const& col,
-//   data_type const output_dtype,
-//   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
-//   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+/**
+ * @brief Compute any of the value in the segment is true when typecasted to bool
+ *
+ * If all elements in input segment are null, the segment result is null.
+ *
+ * @throw cudf::logic_error if input column type is not convertible to bool
+ * @throw cudf::logic_error if `output_dtype` is not bool
+ *
+ * @param col Input column to compute any_of.
+ * @param offsets Indices to identify segment boundaries
+ * @param output_dtype Data type of return type and typecast elements of input column
+ * @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 Column of bool8 for the results of the segments
+ */
+std::unique_ptr<column> segmented_any(
+  column_view const& col,
+  column_view const& offsets,
+  data_type const output_dtype,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
-// /**
-//  * @brief Computes product of elements in input column
-//  *
-//  * If all elements in input column are null, output scalar is null.
-//  *
-//  * @throw cudf::logic_error if input column type is not convertible to `output_dtype`
-//  * @throw cudf::logic_error if `output_dtype` is not arithmetic point type
-//  *
-//  * @param col input column to compute product.
-//  * @param output_dtype data type of return type and typecast elements of input column
-//  * @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 Product as scalar of type `output_dtype`.
-//  */
-// std::unique_ptr<scalar> product(
-//   column_view const& col,
-//   data_type const output_dtype,
-//   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
-//   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+/**
+ * @brief Compute all of the value in the segment is true when typecasted to bool
+ *
+ * If all elements in input segment are null, the segment result is null.
+ *
+ * @throw cudf::logic_error if input column type is not convertible to bool
+ * @throw cudf::logic_error if `output_dtype` is not bool
+ *
+ * @param col Input column to compute all_of.
+ * @param offsets Indices to identify segment boundaries
+ * @param output_dtype Data type of return type and typecast elements of input column
+ * @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 Column of bool8 for the results of the segments
+ */
+std::unique_ptr<column> segmented_all(
+  column_view const& col,
+  column_view const& offsets,
+  data_type const output_dtype,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 }  // namespace reduction
 }  // namespace cudf

cpp/include/cudf/reduction.hpp

@@ -48,6 +48,8 @@ enum class scan_type : bool { INCLUSIVE, EXCLUSIVE };
  * output data type.
  * @throw cudf::logic_error if `min` or `max` reduction is called and the
  * output type does not match the input column data type.
+ * @throw cudf::logic_error if `any` or `all` reduction is called and the
+ * output type is not bool8.
  *
  * If the input column has arithmetic type, output_dtype can be any arithmetic
  * type. For `mean`, `var` and `std` ops, a floating point output type must be
@@ -69,6 +71,41 @@ std::unique_ptr<scalar> reduce(
   data_type output_dtype,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
+/**
+ * @brief  Compute reduction of each segment in the input column
+ *
+ * This function does not detect overflows in reductions.
+ * Using a higher precision `data_type` may prevent overflow.
+ * Only `min` and `max` ops are supported for reduction of non-arithmetic
+ * types (timestamp, string...).
+ * The null values are skipped for the operation.
+ * If the segment is empty, the row corresponding to the result of the
+ * segment is null.
+ *
+ * @throw cudf::logic_error if reduction is called for non-arithmetic output
+ * type and operator other than `min` and `max`.
+ * @throw cudf::logic_error if input column data type is not convertible to
+ * output data type.
+ * @throw cudf::logic_error if `min` or `max` reduction is called and the
+ * output type does not match the input column data type.
+ * @throw cudf::logic_error if `any` or `all` reduction is called and the
+ * output type is not bool8.
+ *
+ * If the input column has arithmetic type, output_dtype can be any arithmetic
+ * type. For `mean`, `var` and `std` ops, a floating point output type must be
+ * specified. If the input column has non-arithmetic type
+ *   eg.(timestamp, string...), the same type must be specified.
+ *
+ * If the reduction fails, the member is_valid of the output scalar
+ * will contain `false`.
+ *
+ * @param col Input column view
+ * @param agg Aggregation operator applied by the reduction
+ * @param offsets Indices to segment boundaries
+ * @param output_dtype  The computation and output precision.
+ * @param mr Device memory resource used to allocate the returned scalar's device memory
+ * @returns Output column with segment's reduce result.
+ */
 std::unique_ptr<column> segmented_reduce(
   column_view const& col,
   column_view const& offsets,

cpp/src/reductions/segmented_all.cu

@@ -0,0 +1,43 @@
+/*
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <cudf/detail/reduction_functions.hpp>
+#include <cudf/detail/utilities/device_atomics.cuh>
+#include <reductions/simple_segmented.cuh>
+
+namespace cudf {
+namespace reduction {
+
+std::unique_ptr<cudf::column> segmented_all(column_view const& col,
+                                            column_view const& offsets,
+                                            cudf::data_type const output_dtype,
+                                            rmm::cuda_stream_view stream,
+                                            rmm::mr::device_memory_resource* mr)
+{
+  CUDF_EXPECTS(output_dtype == cudf::data_type(cudf::type_id::BOOL8),
+               "all() operation can be applied with output type `BOOL8` only");
+
+  // dispatch for non-dictionary types
+  return cudf::type_dispatcher(col.type(),
+                               simple::bool_result_column_dispatcher<cudf::reduction::op::min>{},
+                               col,
+                               offsets,
+                               stream,
+                               mr);
+}
+
+}  // namespace reduction
+}  // namespace cudf