From 90d2cb162785f044e357290779e762ea980f9583 Mon Sep 17 00:00:00 2001 From: David Wendt <45795991+davidwendt@users.noreply.github.com> Date: Thu, 23 Mar 2023 07:00:13 -0400 Subject: [PATCH] Move detail reduction functions to cudf::reduction::detail namespace (#12971) Moves internal/detail reduction (and segmented reduction) functions and other declarations from the `cudf::detail` namespace to the `cudf::reduction::detail` namespace. The detail function headers also moved from `cpp/include/cudf/detail` to `cpp/include/cudf/reduction/detail` directory. The public header `cpp/include/cudf/reduction.hpp` was not changed or moved. Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Vyas Ramasubramani (https://github.com/vyasr) - Robert Maynard (https://github.com/robertmaynard) - AJ Schmidt (https://github.com/ajschmidt8) URL: https://github.com/rapidsai/cudf/pull/12971 --- conda/recipes/libcudf/meta.yaml | 4 +- .../cudf/{ => reduction}/detail/reduction.cuh | 4 +- .../detail/reduction_functions.hpp | 2 + .../detail/reduction_operators.cuh | 9 +- .../detail/segmented_reduction.cuh | 2 +- .../detail/segmented_reduction_functions.hpp | 148 +++++++++--------- .../stream_compaction/apply_boolean_mask.cu | 16 +- cpp/src/reductions/all.cu | 21 +-- cpp/src/reductions/any.cu | 21 +-- cpp/src/reductions/collect_ops.cu | 40 ++--- cpp/src/reductions/compound.cuh | 2 +- cpp/src/reductions/max.cu | 19 ++- cpp/src/reductions/mean.cu | 12 +- cpp/src/reductions/min.cu | 21 ++- cpp/src/reductions/nth_element.cu | 18 ++- cpp/src/reductions/product.cu | 12 +- cpp/src/reductions/reductions.cpp | 103 ++++++------ cpp/src/reductions/segmented/all.cu | 14 +- cpp/src/reductions/segmented/any.cu | 14 +- cpp/src/reductions/segmented/compound.cuh | 2 +- cpp/src/reductions/segmented/max.cu | 15 +- cpp/src/reductions/segmented/mean.cu | 6 +- cpp/src/reductions/segmented/min.cu | 15 +- cpp/src/reductions/segmented/product.cu | 17 +- cpp/src/reductions/segmented/reductions.cpp | 77 +++++---- cpp/src/reductions/segmented/simple.cuh | 16 +- cpp/src/reductions/segmented/std.cu | 7 +- cpp/src/reductions/segmented/sum.cu | 17 +- .../reductions/segmented/sum_of_squares.cu | 6 +- cpp/src/reductions/segmented/var.cu | 7 +- cpp/src/reductions/simple.cuh | 12 +- cpp/src/reductions/std.cu | 10 +- cpp/src/reductions/struct_minmax_util.cuh | 10 +- cpp/src/reductions/sum.cu | 11 +- cpp/src/reductions/sum_of_squares.cu | 11 +- cpp/src/reductions/var.cu | 9 +- 36 files changed, 354 insertions(+), 376 deletions(-) rename cpp/include/cudf/{ => reduction}/detail/reduction.cuh (99%) rename cpp/include/cudf/{ => reduction}/detail/reduction_functions.hpp (99%) rename cpp/include/cudf/{ => reduction}/detail/reduction_operators.cuh (97%) rename cpp/include/cudf/{ => reduction}/detail/segmented_reduction.cuh (99%) rename cpp/include/cudf/{ => reduction}/detail/segmented_reduction_functions.hpp (69%) diff --git a/conda/recipes/libcudf/meta.yaml b/conda/recipes/libcudf/meta.yaml index 0b2fc71aacd..469c25fb673 100644 --- a/conda/recipes/libcudf/meta.yaml +++ b/conda/recipes/libcudf/meta.yaml @@ -114,7 +114,6 @@ outputs: - test -f $PREFIX/include/cudf/detail/nvtx/nvtx3.hpp - test -f $PREFIX/include/cudf/detail/nvtx/ranges.hpp - test -f $PREFIX/include/cudf/detail/quantiles.hpp - - test -f $PREFIX/include/cudf/detail/reduction_functions.hpp - test -f $PREFIX/include/cudf/detail/repeat.hpp - test -f $PREFIX/include/cudf/detail/replace.hpp - test -f $PREFIX/include/cudf/detail/reshape.hpp @@ -123,7 +122,6 @@ outputs: - test -f $PREFIX/include/cudf/detail/scan.hpp - test -f $PREFIX/include/cudf/detail/scatter.hpp - test -f $PREFIX/include/cudf/detail/search.hpp - - test -f $PREFIX/include/cudf/detail/segmented_reduction_functions.hpp - test -f $PREFIX/include/cudf/detail/sequence.hpp - test -f $PREFIX/include/cudf/detail/sorting.hpp - test -f $PREFIX/include/cudf/detail/stream_compaction.hpp @@ -218,6 +216,8 @@ outputs: - test -f $PREFIX/include/cudf/partitioning.hpp - test -f $PREFIX/include/cudf/quantiles.hpp - test -f $PREFIX/include/cudf/reduction.hpp + - test -f $PREFIX/include/cudf/reduction/detail/reduction_functions.hpp + - test -f $PREFIX/include/cudf/reduction/detail/segmented_reduction_functions.hpp - test -f $PREFIX/include/cudf/replace.hpp - test -f $PREFIX/include/cudf/reshape.hpp - test -f $PREFIX/include/cudf/rolling.hpp diff --git a/cpp/include/cudf/detail/reduction.cuh b/cpp/include/cudf/reduction/detail/reduction.cuh similarity index 99% rename from cpp/include/cudf/detail/reduction.cuh rename to cpp/include/cudf/reduction/detail/reduction.cuh index 9dc3b996afc..1620635e0e3 100644 --- a/cpp/include/cudf/detail/reduction.cuh +++ b/cpp/include/cudf/reduction/detail/reduction.cuh @@ -16,7 +16,7 @@ #pragma once -#include +#include "reduction_operators.cuh" #include #include @@ -31,6 +31,8 @@ #include #include +#include + namespace cudf { namespace reduction { namespace detail { diff --git a/cpp/include/cudf/detail/reduction_functions.hpp b/cpp/include/cudf/reduction/detail/reduction_functions.hpp similarity index 99% rename from cpp/include/cudf/detail/reduction_functions.hpp rename to cpp/include/cudf/reduction/detail/reduction_functions.hpp index c554ea6a83e..014a6ba70eb 100644 --- a/cpp/include/cudf/detail/reduction_functions.hpp +++ b/cpp/include/cudf/reduction/detail/reduction_functions.hpp @@ -27,6 +27,7 @@ namespace cudf { namespace reduction { +namespace detail { /** * @brief Computes sum of elements in input column * @@ -323,5 +324,6 @@ std::unique_ptr merge_sets(lists_column_view const& col, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr); +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/include/cudf/detail/reduction_operators.cuh b/cpp/include/cudf/reduction/detail/reduction_operators.cuh similarity index 97% rename from cpp/include/cudf/detail/reduction_operators.cuh rename to cpp/include/cudf/reduction/detail/reduction_operators.cuh index 5a0cb4c1714..0dba84a0b28 100644 --- a/cpp/include/cudf/detail/reduction_operators.cuh +++ b/cpp/include/cudf/reduction/detail/reduction_operators.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -26,6 +26,7 @@ namespace cudf { namespace reduction { +namespace detail { // intermediate data structure to compute `var`, `std` template struct var_std { @@ -244,7 +245,7 @@ struct variance : public compound_op { using op = cudf::DeviceSum; template - using transformer = cudf::reduction::transformer_var_std; + using transformer = cudf::reduction::detail::transformer_var_std; template struct intermediate { @@ -270,7 +271,7 @@ struct standard_deviation : public compound_op { using op = cudf::DeviceSum; template - using transformer = cudf::reduction::transformer_var_std; + using transformer = cudf::reduction::detail::transformer_var_std; template struct intermediate { @@ -288,7 +289,7 @@ struct standard_deviation : public compound_op { }; }; }; - } // namespace op +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/include/cudf/detail/segmented_reduction.cuh b/cpp/include/cudf/reduction/detail/segmented_reduction.cuh similarity index 99% rename from cpp/include/cudf/detail/segmented_reduction.cuh rename to cpp/include/cudf/reduction/detail/segmented_reduction.cuh index 1c39d5eab1e..5c2eaf8cdcb 100644 --- a/cpp/include/cudf/detail/segmented_reduction.cuh +++ b/cpp/include/cudf/reduction/detail/segmented_reduction.cuh @@ -16,7 +16,7 @@ #pragma once -#include +#include "reduction_operators.cuh" #include #include diff --git a/cpp/include/cudf/detail/segmented_reduction_functions.hpp b/cpp/include/cudf/reduction/detail/segmented_reduction_functions.hpp similarity index 69% rename from cpp/include/cudf/detail/segmented_reduction_functions.hpp rename to cpp/include/cudf/reduction/detail/segmented_reduction_functions.hpp index 7b5628fa49a..c1bf59e5f65 100644 --- a/cpp/include/cudf/detail/segmented_reduction_functions.hpp +++ b/cpp/include/cudf/reduction/detail/segmented_reduction_functions.hpp @@ -27,6 +27,7 @@ namespace cudf { namespace reduction { +namespace detail { /** * @brief Compute sum of each segment in the input column @@ -50,14 +51,13 @@ namespace reduction { * @param mr Device memory resource used to allocate the returned column's device memory * @return Sums of segments as type `output_dtype` */ -std::unique_ptr segmented_sum( - column_view const& col, - device_span offsets, - data_type const output_dtype, - null_policy null_handling, - std::optional> init, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr segmented_sum(column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + std::optional> init, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @brief Computes product of each segment in the input column @@ -81,14 +81,13 @@ std::unique_ptr segmented_sum( * @param mr Device memory resource used to allocate the returned column's device memory * @return Product of segments as type `output_dtype` */ -std::unique_ptr segmented_product( - column_view const& col, - device_span offsets, - data_type const output_dtype, - null_policy null_handling, - std::optional> init, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr segmented_product(column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + std::optional> init, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @brief Compute minimum of each segment in the input column @@ -111,14 +110,13 @@ std::unique_ptr segmented_product( * @param mr Device memory resource used to allocate the returned column's device memory * @return Minimums of segments as type `output_dtype` */ -std::unique_ptr segmented_min( - column_view const& col, - device_span offsets, - data_type const output_dtype, - null_policy null_handling, - std::optional> init, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr segmented_min(column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + std::optional> init, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @brief Compute maximum of each segment in the input column @@ -141,14 +139,13 @@ std::unique_ptr segmented_min( * @param mr Device memory resource used to allocate the returned column's device memory * @return Maximums of segments as type `output_dtype` */ -std::unique_ptr segmented_max( - column_view const& col, - device_span offsets, - data_type const output_dtype, - null_policy null_handling, - std::optional> init, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr segmented_max(column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + std::optional> init, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @brief Compute if any of the values in the segment are true when typecasted to bool @@ -172,14 +169,13 @@ std::unique_ptr segmented_max( * @param mr Device memory resource used to allocate the returned column's device memory * @return Column of type BOOL8 for the results of the segments */ -std::unique_ptr segmented_any( - column_view const& col, - device_span offsets, - data_type const output_dtype, - null_policy null_handling, - std::optional> init, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr segmented_any(column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + std::optional> init, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @brief Compute if all of the values in the segment are true when typecasted to bool @@ -203,14 +199,13 @@ std::unique_ptr segmented_any( * @param mr Device memory resource used to allocate the returned column's device memory * @return Column of BOOL8 for the results of the segments */ -std::unique_ptr segmented_all( - column_view const& col, - device_span offsets, - data_type const output_dtype, - null_policy null_handling, - std::optional> init, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr segmented_all(column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + std::optional> init, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @brief Computes mean of elements of segments in the input column @@ -233,13 +228,12 @@ std::unique_ptr segmented_all( * @param mr Device memory resource used to allocate the returned column's device memory * @return Column of `output_dtype` for the reduction results of the segments */ -std::unique_ptr segmented_mean( - column_view const& col, - device_span offsets, - data_type const output_dtype, - null_policy null_handling, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr segmented_mean(column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @brief Computes sum of squares of elements of segments in the input column @@ -262,13 +256,12 @@ std::unique_ptr segmented_mean( * @param mr Device memory resource used to allocate the returned column's device memory * @return Column of `output_dtype` for the reduction results of the segments */ -std::unique_ptr segmented_sum_of_squares( - column_view const& col, - device_span offsets, - data_type const output_dtype, - null_policy null_handling, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr segmented_sum_of_squares(column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @brief Computes the standard deviation of elements of segments in the input column @@ -293,14 +286,13 @@ std::unique_ptr segmented_sum_of_squares( * @param mr Device memory resource used to allocate the returned column's device memory * @return Column of `output_dtype` for the reduction results of the segments */ -std::unique_ptr segmented_standard_deviation( - column_view const& col, - device_span offsets, - data_type const output_dtype, - null_policy null_handling, - size_type ddof, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr segmented_standard_deviation(column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + size_type ddof, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); /** * @brief Computes the variance of elements of segments in the input column @@ -325,14 +317,14 @@ std::unique_ptr segmented_standard_deviation( * @param mr Device memory resource used to allocate the returned column's device memory * @return Column of `output_dtype` for the reduction results of the segments */ -std::unique_ptr segmented_variance( - column_view const& col, - device_span offsets, - data_type const output_dtype, - null_policy null_handling, - size_type ddof, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +std::unique_ptr segmented_variance(column_view const& col, + device_span offsets, + data_type const output_dtype, + null_policy null_handling, + size_type ddof, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/lists/stream_compaction/apply_boolean_mask.cu b/cpp/src/lists/stream_compaction/apply_boolean_mask.cu index 5acb1cb8849..0aaa8356304 100644 --- a/cpp/src/lists/stream_compaction/apply_boolean_mask.cu +++ b/cpp/src/lists/stream_compaction/apply_boolean_mask.cu @@ -21,10 +21,10 @@ #include #include #include -#include #include #include #include +#include #include #include @@ -65,12 +65,14 @@ std::unique_ptr apply_boolean_mask(lists_column_view const& input, cudf::detail::slice( boolean_mask.offsets(), {boolean_mask.offset(), boolean_mask.size() + 1}, stream) .front(); - auto const sizes = cudf::reduction::segmented_sum(boolean_mask_sliced_child, - boolean_mask_sliced_offsets, - offset_data_type, - null_policy::EXCLUDE, - std::nullopt, - stream); + auto const sizes = + cudf::reduction::detail::segmented_sum(boolean_mask_sliced_child, + boolean_mask_sliced_offsets, + offset_data_type, + null_policy::EXCLUDE, + std::nullopt, + stream, + rmm::mr::get_current_device_resource()); auto const d_sizes = column_device_view::create(*sizes, stream); auto const sizes_begin = cudf::detail::make_null_replacement_iterator(*d_sizes, offset_type{0}); auto const sizes_end = sizes_begin + sizes->size(); diff --git a/cpp/src/reductions/all.cu b/cpp/src/reductions/all.cu index 185e14b6e2f..9d32bc4c7f6 100644 --- a/cpp/src/reductions/all.cu +++ b/cpp/src/reductions/all.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -14,10 +14,11 @@ * limitations under the License. */ -#include +#include "simple.cuh" + #include #include -#include +#include #include #include @@ -55,8 +56,7 @@ struct all_fn { { auto const d_dict = cudf::column_device_view::create(input, stream); auto const iter = [&] { - auto null_iter = - cudf::reduction::op::min{}.template get_null_replacing_element_transformer(); + auto null_iter = op::min{}.template get_null_replacing_element_transformer(); auto pair_iter = cudf::dictionary::detail::make_dictionary_pair_iterator(*d_dict, input.has_nulls()); return thrust::make_transform_iterator(pair_iter, null_iter); @@ -78,7 +78,6 @@ struct all_fn { }; } // namespace -} // namespace detail std::unique_ptr all(column_view const& col, cudf::data_type const output_dtype, @@ -93,15 +92,11 @@ std::unique_ptr all(column_view const& col, return cudf::type_dispatcher( dictionary_column_view(col).keys().type(), detail::all_fn{}, col, stream, mr); } + using reducer = simple::detail::bool_result_element_dispatcher; // dispatch for non-dictionary types - return cudf::type_dispatcher( - col.type(), - simple::detail::bool_result_element_dispatcher{}, - col, - init, - stream, - mr); + return cudf::type_dispatcher(col.type(), reducer{}, col, init, stream, mr); } +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/any.cu b/cpp/src/reductions/any.cu index 871672e5c03..07977d2417f 100644 --- a/cpp/src/reductions/any.cu +++ b/cpp/src/reductions/any.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -14,10 +14,11 @@ * limitations under the License. */ -#include +#include "simple.cuh" + #include #include -#include +#include #include #include @@ -55,8 +56,7 @@ struct any_fn { { auto const d_dict = cudf::column_device_view::create(input, stream); auto const iter = [&] { - auto null_iter = - cudf::reduction::op::max{}.template get_null_replacing_element_transformer(); + auto null_iter = op::max{}.template get_null_replacing_element_transformer(); auto pair_iter = cudf::dictionary::detail::make_dictionary_pair_iterator(*d_dict, input.has_nulls()); return thrust::make_transform_iterator(pair_iter, null_iter); @@ -78,7 +78,6 @@ struct any_fn { }; } // namespace -} // namespace detail std::unique_ptr any(column_view const& col, cudf::data_type const output_dtype, @@ -93,15 +92,11 @@ std::unique_ptr any(column_view const& col, return cudf::type_dispatcher( dictionary_column_view(col).keys().type(), detail::any_fn{}, col, stream, mr); } + using reducer = simple::detail::bool_result_element_dispatcher; // dispatch for non-dictionary types - return cudf::type_dispatcher( - col.type(), - simple::detail::bool_result_element_dispatcher{}, - col, - init, - stream, - mr); + return cudf::type_dispatcher(col.type(), reducer{}, col, init, stream, mr); } +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/collect_ops.cu b/cpp/src/reductions/collect_ops.cu index 4d6a32b528a..743eddbffaf 100644 --- a/cpp/src/reductions/collect_ops.cu +++ b/cpp/src/reductions/collect_ops.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -17,15 +17,15 @@ #include #include #include -#include #include #include +#include #include #include namespace cudf { namespace reduction { - +namespace detail { namespace { /** @@ -49,8 +49,8 @@ std::unique_ptr collect_list(column_view const& col, { if (need_handle_nulls(col, null_handling)) { auto d_view = column_device_view::create(col, stream); - auto filter = detail::validity_accessor(*d_view); - auto null_purged_table = detail::copy_if(table_view{{col}}, filter, stream, mr); + auto filter = cudf::detail::validity_accessor(*d_view); + auto null_purged_table = cudf::detail::copy_if(table_view{{col}}, filter, stream, mr); column* null_purged_col = null_purged_table->release().front().release(); null_purged_col->set_null_mask(rmm::device_buffer{0, stream, mr}, 0); return std::make_unique(std::move(*null_purged_col), true, stream, mr); @@ -86,13 +86,13 @@ std::unique_ptr collect_set(column_view const& col, return std::pair(col, std::unique_ptr(nullptr)); }(); - auto distinct_table = detail::distinct(table_view{{input_as_collect_list}}, - std::vector{0}, - duplicate_keep_option::KEEP_ANY, - nulls_equal, - nans_equal, - stream, - mr); + auto distinct_table = cudf::detail::distinct(table_view{{input_as_collect_list}}, + std::vector{0}, + duplicate_keep_option::KEEP_ANY, + nulls_equal, + nans_equal, + stream, + mr); return std::make_unique(std::move(distinct_table->get_column(0)), true, stream, mr); } @@ -104,15 +104,15 @@ std::unique_ptr merge_sets(lists_column_view const& col, rmm::mr::device_memory_resource* mr) { auto flatten_col = col.get_sliced_child(stream); - auto distinct_table = detail::distinct(table_view{{flatten_col}}, - std::vector{0}, - duplicate_keep_option::KEEP_ANY, - nulls_equal, - nans_equal, - stream, - mr); + auto distinct_table = cudf::detail::distinct(table_view{{flatten_col}}, + std::vector{0}, + duplicate_keep_option::KEEP_ANY, + nulls_equal, + nans_equal, + stream, + mr); return std::make_unique(std::move(distinct_table->get_column(0)), true, stream, mr); } - +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/compound.cuh b/cpp/src/reductions/compound.cuh index 9458ae2d581..3428130d912 100644 --- a/cpp/src/reductions/compound.cuh +++ b/cpp/src/reductions/compound.cuh @@ -16,8 +16,8 @@ #pragma once -#include #include +#include #include #include #include diff --git a/cpp/src/reductions/max.cu b/cpp/src/reductions/max.cu index b57896e5fc0..1cf2b6f53b6 100644 --- a/cpp/src/reductions/max.cu +++ b/cpp/src/reductions/max.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -14,14 +14,16 @@ * limitations under the License. */ -#include +#include "simple.cuh" + #include -#include +#include #include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr max(column_view const& col, cudf::data_type const output_dtype, @@ -35,14 +37,11 @@ std::unique_ptr 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::detail::same_element_type_dispatcher{}, - col, - init, - stream, - mr); + + using reducer = simple::detail::same_element_type_dispatcher; + return cudf::type_dispatcher(dispatch_type, reducer{}, col, init, stream, mr); } +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/mean.cu b/cpp/src/reductions/mean.cu index e4b5f754b9b..e64660932ce 100644 --- a/cpp/src/reductions/mean.cu +++ b/cpp/src/reductions/mean.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -14,26 +14,30 @@ * limitations under the License. */ -#include +#include "compound.cuh" + #include -#include +#include #include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr mean(column_view const& col, cudf::data_type const output_dtype, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - using reducer = compound::detail::element_type_dispatcher; auto col_type = cudf::is_dictionary(col.type()) ? dictionary_column_view(col).keys().type() : col.type(); + + using reducer = compound::detail::element_type_dispatcher; return cudf::type_dispatcher( col_type, reducer(), col, output_dtype, /* ddof is not used for mean*/ 1, stream, mr); } +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/min.cu b/cpp/src/reductions/min.cu index ed16cec5ffd..792965e8b99 100644 --- a/cpp/src/reductions/min.cu +++ b/cpp/src/reductions/min.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -14,13 +14,14 @@ * limitations under the License. */ -#include +#include "simple.cuh" + #include -#include +#include namespace cudf { namespace reduction { - +namespace detail { std::unique_ptr min(column_view const& col, data_type const output_dtype, std::optional> init, @@ -33,14 +34,10 @@ std::unique_ptr 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::detail::same_element_type_dispatcher{}, - col, - init, - stream, - mr); -} + using reducer = simple::detail::same_element_type_dispatcher; + return cudf::type_dispatcher(dispatch_type, reducer{}, col, init, stream, mr); +} +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/nth_element.cu b/cpp/src/reductions/nth_element.cu index 78c469ee767..ef58ec3f42e 100644 --- a/cpp/src/reductions/nth_element.cu +++ b/cpp/src/reductions/nth_element.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -18,7 +18,7 @@ #include #include #include -#include +#include #include #include @@ -28,11 +28,13 @@ #include #include -std::unique_ptr cudf::reduction::nth_element(column_view const& col, - size_type n, - null_policy null_handling, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +namespace cudf::reduction::detail { + +std::unique_ptr nth_element(column_view const& col, + size_type n, + null_policy null_handling, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { CUDF_EXPECTS(n >= -col.size() and n < col.size(), "Index out of bounds"); auto wrap_n = [n](size_type size) { return (n < 0 ? size + n : n); }; @@ -60,3 +62,5 @@ std::unique_ptr cudf::reduction::nth_element(column_view const& co return cudf::detail::get_element(col, n, stream, mr); } } + +} // namespace cudf::reduction::detail diff --git a/cpp/src/reductions/product.cu b/cpp/src/reductions/product.cu index 39e031f69d1..2e483813939 100644 --- a/cpp/src/reductions/product.cu +++ b/cpp/src/reductions/product.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -14,14 +14,16 @@ * limitations under the License. */ -#include +#include "simple.cuh" + #include -#include +#include #include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr product(column_view const& col, cudf::data_type const output_dtype, @@ -31,13 +33,13 @@ std::unique_ptr product(column_view const& col, { return cudf::type_dispatcher( cudf::is_dictionary(col.type()) ? dictionary_column_view(col).keys().type() : col.type(), - simple::detail::element_type_dispatcher{}, + simple::detail::element_type_dispatcher{}, col, output_dtype, init, stream, mr); } - +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/reductions.cpp b/cpp/src/reductions/reductions.cpp index cae2699aac7..b6c050287cf 100644 --- a/cpp/src/reductions/reductions.cpp +++ b/cpp/src/reductions/reductions.cpp @@ -19,11 +19,11 @@ #include #include #include -#include #include #include #include #include +#include #include #include #include @@ -31,6 +31,7 @@ #include namespace cudf { +namespace reduction { namespace detail { struct reduce_dispatch_functor { column_view const col; @@ -52,89 +53,89 @@ struct reduce_dispatch_functor { std::unique_ptr operator()(reduce_aggregation const& agg) { switch (k) { - case aggregation::SUM: return reduction::sum(col, output_dtype, init, stream, mr); - case aggregation::PRODUCT: return reduction::product(col, output_dtype, init, stream, mr); - case aggregation::MIN: return reduction::min(col, output_dtype, init, stream, mr); - case aggregation::MAX: return reduction::max(col, output_dtype, init, stream, mr); - case aggregation::ANY: return reduction::any(col, output_dtype, init, stream, mr); - case aggregation::ALL: return reduction::all(col, output_dtype, init, stream, mr); - case aggregation::SUM_OF_SQUARES: - return reduction::sum_of_squares(col, output_dtype, stream, mr); - case aggregation::MEAN: return reduction::mean(col, output_dtype, stream, mr); + case aggregation::SUM: return sum(col, output_dtype, init, stream, mr); + case aggregation::PRODUCT: return product(col, output_dtype, init, stream, mr); + case aggregation::MIN: return min(col, output_dtype, init, stream, mr); + case aggregation::MAX: return max(col, output_dtype, init, stream, mr); + case aggregation::ANY: return any(col, output_dtype, init, stream, mr); + case aggregation::ALL: return all(col, output_dtype, init, stream, mr); + case aggregation::SUM_OF_SQUARES: return sum_of_squares(col, output_dtype, stream, mr); + case aggregation::MEAN: return mean(col, output_dtype, stream, mr); case aggregation::VARIANCE: { - auto var_agg = static_cast(agg); - return reduction::variance(col, output_dtype, var_agg._ddof, stream, mr); + auto var_agg = static_cast(agg); + return variance(col, output_dtype, var_agg._ddof, stream, mr); } case aggregation::STD: { - auto var_agg = static_cast(agg); - return reduction::standard_deviation(col, output_dtype, var_agg._ddof, stream, mr); + auto var_agg = static_cast(agg); + return standard_deviation(col, output_dtype, var_agg._ddof, stream, mr); } case aggregation::MEDIAN: { - auto current_mr = rmm::mr::get_current_device_resource(); - auto sorted_indices = - sorted_order(table_view{{col}}, {}, {null_order::AFTER}, stream, current_mr); + auto current_mr = rmm::mr::get_current_device_resource(); + auto sorted_indices = cudf::detail::sorted_order( + table_view{{col}}, {}, {null_order::AFTER}, stream, current_mr); auto valid_sorted_indices = - split(*sorted_indices, {col.size() - col.null_count()}, stream)[0]; - auto col_ptr = quantile( + cudf::detail::split(*sorted_indices, {col.size() - col.null_count()}, stream)[0]; + auto col_ptr = cudf::detail::quantile( col, {0.5}, interpolation::LINEAR, valid_sorted_indices, true, stream, current_mr); - return get_element(*col_ptr, 0, stream, mr); + return cudf::detail::get_element(*col_ptr, 0, stream, mr); } case aggregation::QUANTILE: { - auto quantile_agg = static_cast(agg); + auto quantile_agg = static_cast(agg); CUDF_EXPECTS(quantile_agg._quantiles.size() == 1, "Reduction quantile accepts only one quantile value"); - auto current_mr = rmm::mr::get_current_device_resource(); - auto sorted_indices = - sorted_order(table_view{{col}}, {}, {null_order::AFTER}, stream, current_mr); + auto current_mr = rmm::mr::get_current_device_resource(); + auto sorted_indices = cudf::detail::sorted_order( + table_view{{col}}, {}, {null_order::AFTER}, stream, current_mr); auto valid_sorted_indices = - split(*sorted_indices, {col.size() - col.null_count()}, stream)[0]; + cudf::detail::split(*sorted_indices, {col.size() - col.null_count()}, stream)[0]; - auto col_ptr = quantile(col, - quantile_agg._quantiles, - quantile_agg._interpolation, - valid_sorted_indices, - true, - stream, - current_mr); - return get_element(*col_ptr, 0, stream, mr); + auto col_ptr = cudf::detail::quantile(col, + quantile_agg._quantiles, + quantile_agg._interpolation, + valid_sorted_indices, + true, + stream, + current_mr); + return cudf::detail::get_element(*col_ptr, 0, stream, mr); } case aggregation::NUNIQUE: { - auto nunique_agg = static_cast(agg); - return make_fixed_width_scalar( - detail::distinct_count(col, nunique_agg._null_handling, nan_policy::NAN_IS_VALID, stream), + auto nunique_agg = static_cast(agg); + return cudf::make_fixed_width_scalar( + cudf::detail::distinct_count( + col, nunique_agg._null_handling, nan_policy::NAN_IS_VALID, stream), stream, mr); } case aggregation::NTH_ELEMENT: { - auto nth_agg = static_cast(agg); - return reduction::nth_element(col, nth_agg._n, nth_agg._null_handling, stream, mr); + auto nth_agg = static_cast(agg); + return nth_element(col, nth_agg._n, nth_agg._null_handling, stream, mr); } case aggregation::COLLECT_LIST: { - auto col_agg = static_cast(agg); - return reduction::collect_list(col, col_agg._null_handling, stream, mr); + auto col_agg = static_cast(agg); + return collect_list(col, col_agg._null_handling, stream, mr); } case aggregation::COLLECT_SET: { - auto col_agg = static_cast(agg); - return reduction::collect_set( + auto col_agg = static_cast(agg); + return collect_set( col, col_agg._null_handling, col_agg._nulls_equal, col_agg._nans_equal, stream, mr); } case aggregation::MERGE_LISTS: { - return reduction::merge_lists(col, stream, mr); + return merge_lists(col, stream, mr); } case aggregation::MERGE_SETS: { - auto col_agg = static_cast(agg); - return reduction::merge_sets(col, col_agg._nulls_equal, col_agg._nans_equal, stream, mr); + auto col_agg = static_cast(agg); + return merge_sets(col, col_agg._nulls_equal, col_agg._nans_equal, stream, mr); } case aggregation::TDIGEST: { CUDF_EXPECTS(output_dtype.id() == type_id::STRUCT, "Tdigest aggregations expect output type to be STRUCT"); - auto td_agg = static_cast(agg); + auto td_agg = static_cast(agg); return tdigest::detail::reduce_tdigest(col, td_agg.max_centroids, stream, mr); } case aggregation::MERGE_TDIGEST: { CUDF_EXPECTS(output_dtype.id() == type_id::STRUCT, "Tdigest aggregations expect output type to be STRUCT"); - auto td_agg = static_cast(agg); + auto td_agg = static_cast(agg); return tdigest::detail::reduce_merge_tdigest(col, td_agg.max_centroids, stream, mr); } default: CUDF_FAIL("Unsupported reduction operator"); @@ -183,10 +184,11 @@ std::unique_ptr reduce(column_view const& col, return result; } - return aggregation_dispatcher( + return cudf::detail::aggregation_dispatcher( agg.kind, reduce_dispatch_functor{col, output_dtype, init, stream, mr}, agg); } } // namespace detail +} // namespace reduction std::unique_ptr reduce(column_view const& col, reduce_aggregation const& agg, @@ -194,7 +196,8 @@ std::unique_ptr reduce(column_view const& col, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::reduce(col, agg, output_dtype, std::nullopt, cudf::get_default_stream(), mr); + return reduction::detail::reduce( + col, agg, output_dtype, std::nullopt, cudf::get_default_stream(), mr); } std::unique_ptr reduce(column_view const& col, @@ -204,6 +207,6 @@ std::unique_ptr reduce(column_view const& col, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::reduce(col, agg, output_dtype, init, cudf::get_default_stream(), mr); + return reduction::detail::reduce(col, agg, output_dtype, init, cudf::get_default_stream(), mr); } } // namespace cudf diff --git a/cpp/src/reductions/segmented/all.cu b/cpp/src/reductions/segmented/all.cu index f75fcd8066c..b81a088155c 100644 --- a/cpp/src/reductions/segmented/all.cu +++ b/cpp/src/reductions/segmented/all.cu @@ -16,10 +16,11 @@ #include "simple.cuh" -#include +#include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr segmented_all( column_view const& col, @@ -33,17 +34,12 @@ std::unique_ptr segmented_all( CUDF_EXPECTS(output_dtype == cudf::data_type(cudf::type_id::BOOL8), "segmented_all() operation requires output type `BOOL8`"); + using reducer = simple::detail::bool_result_column_dispatcher; // A minimum over bool types is used to implement all() return cudf::type_dispatcher( - col.type(), - simple::detail::bool_result_column_dispatcher{}, - col, - offsets, - null_handling, - init, - stream, - mr); + col.type(), reducer{}, col, offsets, null_handling, init, stream, mr); } +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/segmented/any.cu b/cpp/src/reductions/segmented/any.cu index 6a4fc70d438..9210fbd3c7c 100644 --- a/cpp/src/reductions/segmented/any.cu +++ b/cpp/src/reductions/segmented/any.cu @@ -16,10 +16,11 @@ #include "simple.cuh" -#include +#include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr segmented_any( column_view const& col, @@ -33,17 +34,12 @@ std::unique_ptr segmented_any( CUDF_EXPECTS(output_dtype == cudf::data_type(cudf::type_id::BOOL8), "segmented_any() operation requires output type `BOOL8`"); + using reducer = simple::detail::bool_result_column_dispatcher; // A maximum over bool types is used to implement any() return cudf::type_dispatcher( - col.type(), - simple::detail::bool_result_column_dispatcher{}, - col, - offsets, - null_handling, - init, - stream, - mr); + col.type(), reducer{}, col, offsets, null_handling, init, stream, mr); } +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/segmented/compound.cuh b/cpp/src/reductions/segmented/compound.cuh index e8abd32cf61..395ad4c1dc9 100644 --- a/cpp/src/reductions/segmented/compound.cuh +++ b/cpp/src/reductions/segmented/compound.cuh @@ -21,7 +21,7 @@ #include #include -#include +#include #include #include diff --git a/cpp/src/reductions/segmented/max.cu b/cpp/src/reductions/segmented/max.cu index d72b65301c1..c07c8fb2269 100644 --- a/cpp/src/reductions/segmented/max.cu +++ b/cpp/src/reductions/segmented/max.cu @@ -16,10 +16,11 @@ #include "simple.cuh" -#include +#include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr segmented_max( column_view const& col, @@ -32,16 +33,10 @@ std::unique_ptr segmented_max( { CUDF_EXPECTS(col.type() == output_dtype, "segmented_max() operation requires matching output type"); + using reducer = simple::detail::same_column_type_dispatcher; return cudf::type_dispatcher( - col.type(), - simple::detail::same_column_type_dispatcher{}, - col, - offsets, - null_handling, - init, - stream, - mr); + col.type(), reducer{}, col, offsets, null_handling, init, stream, mr); } - +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/segmented/mean.cu b/cpp/src/reductions/segmented/mean.cu index b7a5bfa43d6..99f1533a154 100644 --- a/cpp/src/reductions/segmented/mean.cu +++ b/cpp/src/reductions/segmented/mean.cu @@ -16,12 +16,13 @@ #include "compound.cuh" -#include +#include #include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr segmented_mean(column_view const& col, device_span offsets, @@ -30,11 +31,12 @@ std::unique_ptr segmented_mean(column_view const& col, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - using reducer = compound::detail::compound_segmented_dispatcher; + using reducer = compound::detail::compound_segmented_dispatcher; constexpr size_type ddof = 1; // ddof for mean calculation return cudf::type_dispatcher( col.type(), reducer{}, col, offsets, output_dtype, null_handling, ddof, stream, mr); } +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/segmented/min.cu b/cpp/src/reductions/segmented/min.cu index b7fbedf2690..f1597f90267 100644 --- a/cpp/src/reductions/segmented/min.cu +++ b/cpp/src/reductions/segmented/min.cu @@ -16,10 +16,11 @@ #include "simple.cuh" -#include +#include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr segmented_min( column_view const& col, @@ -32,16 +33,10 @@ std::unique_ptr segmented_min( { CUDF_EXPECTS(col.type() == output_dtype, "segmented_min() operation requires matching output type"); + using reducer = simple::detail::same_column_type_dispatcher; return cudf::type_dispatcher( - col.type(), - simple::detail::same_column_type_dispatcher{}, - col, - offsets, - null_handling, - init, - stream, - mr); + col.type(), reducer{}, col, offsets, null_handling, init, stream, mr); } - +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/segmented/product.cu b/cpp/src/reductions/segmented/product.cu index d5442126660..ea9c6f484c0 100644 --- a/cpp/src/reductions/segmented/product.cu +++ b/cpp/src/reductions/segmented/product.cu @@ -16,11 +16,11 @@ #include "simple.cuh" -#include +#include namespace cudf { namespace reduction { - +namespace detail { std::unique_ptr segmented_product( column_view const& col, device_span offsets, @@ -30,17 +30,10 @@ std::unique_ptr segmented_product( rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { + using reducer = simple::detail::column_type_dispatcher; return cudf::type_dispatcher( - col.type(), - simple::detail::column_type_dispatcher{}, - col, - offsets, - output_dtype, - null_handling, - init, - stream, - mr); + col.type(), reducer{}, col, offsets, output_dtype, null_handling, init, stream, mr); } - +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/segmented/reductions.cpp b/cpp/src/reductions/segmented/reductions.cpp index 1de55b371b3..66b98fa8322 100644 --- a/cpp/src/reductions/segmented/reductions.cpp +++ b/cpp/src/reductions/segmented/reductions.cpp @@ -18,7 +18,7 @@ #include #include #include -#include +#include #include #include #include @@ -26,6 +26,7 @@ #include namespace cudf { +namespace reduction { namespace detail { struct segmented_reduce_dispatch_functor { column_view const& col; @@ -69,41 +70,32 @@ struct segmented_reduce_dispatch_functor { { switch (k) { case segmented_reduce_aggregation::SUM: - return reduction::segmented_sum( - col, offsets, output_dtype, null_handling, init, stream, mr); + return segmented_sum(col, offsets, output_dtype, null_handling, init, stream, mr); case segmented_reduce_aggregation::PRODUCT: - return reduction::segmented_product( - col, offsets, output_dtype, null_handling, init, stream, mr); + return segmented_product(col, offsets, output_dtype, null_handling, init, stream, mr); case segmented_reduce_aggregation::MIN: - return reduction::segmented_min( - col, offsets, output_dtype, null_handling, init, stream, mr); + return segmented_min(col, offsets, output_dtype, null_handling, init, stream, mr); case segmented_reduce_aggregation::MAX: - return reduction::segmented_max( - col, offsets, output_dtype, null_handling, init, stream, mr); + return segmented_max(col, offsets, output_dtype, null_handling, init, stream, mr); case segmented_reduce_aggregation::ANY: - return reduction::segmented_any( - col, offsets, output_dtype, null_handling, init, stream, mr); + return segmented_any(col, offsets, output_dtype, null_handling, init, stream, mr); case segmented_reduce_aggregation::ALL: - return reduction::segmented_all( - col, offsets, output_dtype, null_handling, init, stream, mr); + return segmented_all(col, offsets, output_dtype, null_handling, init, stream, mr); case segmented_reduce_aggregation::SUM_OF_SQUARES: - return reduction::segmented_sum_of_squares( - col, offsets, output_dtype, null_handling, stream, mr); + return segmented_sum_of_squares(col, offsets, output_dtype, null_handling, stream, mr); case segmented_reduce_aggregation::MEAN: - return reduction::segmented_mean(col, offsets, output_dtype, null_handling, stream, mr); - case aggregation::VARIANCE: { - auto var_agg = static_cast(agg); - return reduction::segmented_variance( + return segmented_mean(col, offsets, output_dtype, null_handling, stream, mr); + case segmented_reduce_aggregation::VARIANCE: { + auto var_agg = static_cast(agg); + return segmented_variance( col, offsets, output_dtype, null_handling, var_agg._ddof, stream, mr); } - case aggregation::STD: { - auto var_agg = static_cast(agg); - return reduction::segmented_standard_deviation( + case segmented_reduce_aggregation::STD: { + auto var_agg = static_cast(agg); + return segmented_standard_deviation( col, offsets, output_dtype, null_handling, var_agg._ddof, stream, mr); } - default: - CUDF_FAIL("Unsupported aggregation type."); - // TODO: Add support for compound_ops. GH #10432 + default: CUDF_FAIL("Unsupported aggregation type."); } } }; @@ -127,13 +119,14 @@ std::unique_ptr segmented_reduce(column_view const& segmented_values, } CUDF_EXPECTS(offsets.size() > 0, "`offsets` should have at least 1 element."); - return aggregation_dispatcher( + return cudf::detail::aggregation_dispatcher( agg.kind, segmented_reduce_dispatch_functor{ segmented_values, offsets, output_dtype, null_handling, init, stream, mr}, agg); } } // namespace detail +} // namespace reduction std::unique_ptr segmented_reduce(column_view const& segmented_values, device_span offsets, @@ -143,14 +136,14 @@ std::unique_ptr segmented_reduce(column_view const& segmented_values, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::segmented_reduce(segmented_values, - offsets, - agg, - output_dtype, - null_handling, - std::nullopt, - cudf::get_default_stream(), - mr); + return reduction::detail::segmented_reduce(segmented_values, + offsets, + agg, + output_dtype, + null_handling, + std::nullopt, + cudf::get_default_stream(), + mr); } std::unique_ptr segmented_reduce(column_view const& segmented_values, @@ -162,14 +155,14 @@ std::unique_ptr segmented_reduce(column_view const& segmented_values, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::segmented_reduce(segmented_values, - offsets, - agg, - output_dtype, - null_handling, - init, - cudf::get_default_stream(), - mr); + return reduction::detail::segmented_reduce(segmented_values, + offsets, + agg, + output_dtype, + null_handling, + init, + cudf::get_default_stream(), + mr); } } // namespace cudf diff --git a/cpp/src/reductions/segmented/simple.cuh b/cpp/src/reductions/segmented/simple.cuh index 0c22848fd89..32138f0835b 100644 --- a/cpp/src/reductions/segmented/simple.cuh +++ b/cpp/src/reductions/segmented/simple.cuh @@ -22,11 +22,11 @@ #include #include #include -#include #include #include #include #include +#include #include #include #include @@ -133,8 +133,8 @@ std::unique_ptr simple_segmented_reduction( template || - std::is_same_v)> + CUDF_ENABLE_IF(std::is_same_v || + std::is_same_v)> std::unique_ptr string_segmented_reduction(column_view const& col, device_span offsets, null_policy null_handling, @@ -147,7 +147,7 @@ std::unique_ptr string_segmented_reduction(column_view const& col, auto it = thrust::make_counting_iterator(0); auto const num_segments = static_cast(offsets.size()) - 1; - bool constexpr is_argmin = std::is_same_v; + bool constexpr is_argmin = std::is_same_v; auto string_comparator = cudf::detail::element_argminmax_fn{*device_col, col.has_nulls(), is_argmin}; auto constexpr identity = @@ -178,8 +178,8 @@ std::unique_ptr string_segmented_reduction(column_view const& col, template () && - !std::is_same_v())> + CUDF_ENABLE_IF(!std::is_same_v() && + !std::is_same_v())> std::unique_ptr string_segmented_reduction(column_view const& col, device_span offsets, null_policy null_handling, @@ -215,7 +215,7 @@ std::unique_ptr fixed_point_segmented_reduction( auto result = simple_segmented_reduction(col, offsets, null_handling, init, stream, mr); auto const scale = [&] { - if constexpr (std::is_same_v) { + if constexpr (std::is_same_v) { // The product aggregation requires updating the scale of the fixed-point output column. // The output scale needs to be the maximum count of all segments multiplied by // the input scale value. @@ -245,7 +245,7 @@ std::unique_ptr fixed_point_segmented_reduction( return new_scale; } - if constexpr (std::is_same_v) { + if constexpr (std::is_same_v) { return numeric::scale_type{col.type().scale() * 2}; } diff --git a/cpp/src/reductions/segmented/std.cu b/cpp/src/reductions/segmented/std.cu index 6af5a9cf9b6..5f5ced63b8f 100644 --- a/cpp/src/reductions/segmented/std.cu +++ b/cpp/src/reductions/segmented/std.cu @@ -16,12 +16,13 @@ #include "compound.cuh" -#include +#include #include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr segmented_standard_deviation(column_view const& col, device_span offsets, @@ -31,11 +32,11 @@ std::unique_ptr segmented_standard_deviation(column_view const& co rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - using reducer = - compound::detail::compound_segmented_dispatcher; + using reducer = compound::detail::compound_segmented_dispatcher; return cudf::type_dispatcher( col.type(), reducer(), col, offsets, output_dtype, null_handling, ddof, stream, mr); } +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/segmented/sum.cu b/cpp/src/reductions/segmented/sum.cu index 0cb8decdc58..7e84961dee0 100644 --- a/cpp/src/reductions/segmented/sum.cu +++ b/cpp/src/reductions/segmented/sum.cu @@ -16,10 +16,11 @@ #include "simple.cuh" -#include +#include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr segmented_sum( column_view const& col, @@ -30,16 +31,10 @@ std::unique_ptr segmented_sum( rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - return cudf::type_dispatcher(col.type(), - simple::detail::column_type_dispatcher{}, - col, - offsets, - output_dtype, - null_handling, - init, - stream, - mr); + using reducer = simple::detail::column_type_dispatcher; + return cudf::type_dispatcher( + col.type(), reducer{}, col, offsets, output_dtype, null_handling, init, stream, mr); } - +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/segmented/sum_of_squares.cu b/cpp/src/reductions/segmented/sum_of_squares.cu index 1ee4f992b6d..6c3f286fd8d 100644 --- a/cpp/src/reductions/segmented/sum_of_squares.cu +++ b/cpp/src/reductions/segmented/sum_of_squares.cu @@ -16,12 +16,13 @@ #include "simple.cuh" -#include +#include #include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr segmented_sum_of_squares(column_view const& col, device_span offsets, @@ -30,10 +31,11 @@ std::unique_ptr segmented_sum_of_squares(column_view const& col, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - using reducer = simple::detail::column_type_dispatcher; + using reducer = simple::detail::column_type_dispatcher; return cudf::type_dispatcher( col.type(), reducer{}, col, offsets, output_dtype, null_handling, std::nullopt, stream, mr); } +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/segmented/var.cu b/cpp/src/reductions/segmented/var.cu index 84adf353700..4ac815b542f 100644 --- a/cpp/src/reductions/segmented/var.cu +++ b/cpp/src/reductions/segmented/var.cu @@ -16,12 +16,13 @@ #include "compound.cuh" -#include +#include #include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr segmented_variance(column_view const& col, device_span offsets, @@ -31,10 +32,10 @@ std::unique_ptr segmented_variance(column_view const& col, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - using reducer = compound::detail::compound_segmented_dispatcher; + using reducer = compound::detail::compound_segmented_dispatcher; return cudf::type_dispatcher( col.type(), reducer(), col, offsets, output_dtype, null_handling, ddof, stream, mr); } - +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/simple.cuh b/cpp/src/reductions/simple.cuh index c7c0d400106..189c17f9b28 100644 --- a/cpp/src/reductions/simple.cuh +++ b/cpp/src/reductions/simple.cuh @@ -16,13 +16,13 @@ #pragma once -#include +#include "struct_minmax_util.cuh" #include -#include #include #include #include +#include #include #include #include @@ -117,10 +117,10 @@ std::unique_ptr fixed_point_reduction( auto result = simple_reduction(col, init, stream, mr); auto const scale = [&] { - if (std::is_same_v) { + if (std::is_same_v) { auto const valid_count = static_cast(col.size() - col.null_count()); return numeric::scale_type{col.type().scale() * (valid_count + (init.has_value() ? 1 : 0))}; - } else if (std::is_same_v) { + } else if (std::is_same_v) { return numeric::scale_type{col.type().scale() * 2}; } return numeric::scale_type{col.type().scale()}; @@ -300,8 +300,8 @@ struct same_element_type_dispatcher { public: template && - (std::is_same_v || - std::is_same_v)>* = nullptr> + (std::is_same_v || + std::is_same_v)>* = nullptr> std::unique_ptr operator()(column_view const& input, std::optional> init, rmm::cuda_stream_view stream, diff --git a/cpp/src/reductions/std.cu b/cpp/src/reductions/std.cu index e9ba75f68e6..9df83634667 100644 --- a/cpp/src/reductions/std.cu +++ b/cpp/src/reductions/std.cu @@ -14,14 +14,16 @@ * limitations under the License. */ -#include +#include "compound.cuh" + #include -#include +#include #include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr standard_deviation(column_view const& col, cudf::data_type const output_dtype, @@ -31,8 +33,7 @@ std::unique_ptr standard_deviation(column_view const& col, { // TODO: add cuda version check when the fix is available #if !defined(__CUDACC_DEBUG__) - using reducer = - compound::detail::element_type_dispatcher; + using reducer = compound::detail::element_type_dispatcher; 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); @@ -43,5 +44,6 @@ std::unique_ptr standard_deviation(column_view const& col, #endif } +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/struct_minmax_util.cuh b/cpp/src/reductions/struct_minmax_util.cuh index b2106066ff2..f8f7ee84e34 100644 --- a/cpp/src/reductions/struct_minmax_util.cuh +++ b/cpp/src/reductions/struct_minmax_util.cuh @@ -17,10 +17,10 @@ #pragma once #include -#include #include #include #include +#include #include #include #include @@ -134,10 +134,10 @@ class comparison_binop_generator { template static auto create(column_view const& input, rmm::cuda_stream_view stream) { - return comparison_binop_generator( - input, - stream, - std::is_same_v || std::is_same_v); + return comparison_binop_generator(input, + stream, + std::is_same_v || + std::is_same_v); } template diff --git a/cpp/src/reductions/sum.cu b/cpp/src/reductions/sum.cu index b919d871cc2..85c6b32dbaf 100644 --- a/cpp/src/reductions/sum.cu +++ b/cpp/src/reductions/sum.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -14,14 +14,16 @@ * limitations under the License. */ -#include +#include "simple.cuh" + #include -#include +#include #include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr sum(column_view const& col, cudf::data_type const output_dtype, @@ -31,7 +33,7 @@ std::unique_ptr sum(column_view const& col, { return cudf::type_dispatcher( cudf::is_dictionary(col.type()) ? dictionary_column_view(col).keys().type() : col.type(), - simple::detail::element_type_dispatcher{}, + simple::detail::element_type_dispatcher{}, col, output_dtype, init, @@ -39,5 +41,6 @@ std::unique_ptr sum(column_view const& col, mr); } +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/sum_of_squares.cu b/cpp/src/reductions/sum_of_squares.cu index af28ba19c9a..7b85c4e6dc9 100644 --- a/cpp/src/reductions/sum_of_squares.cu +++ b/cpp/src/reductions/sum_of_squares.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -14,14 +14,16 @@ * limitations under the License. */ -#include +#include "simple.cuh" + #include -#include +#include #include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr sum_of_squares(column_view const& col, cudf::data_type const output_dtype, @@ -30,7 +32,7 @@ std::unique_ptr 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::detail::element_type_dispatcher{}, + simple::detail::element_type_dispatcher{}, col, output_dtype, std::nullopt, @@ -38,5 +40,6 @@ std::unique_ptr sum_of_squares(column_view const& col, mr); } +} // namespace detail } // namespace reduction } // namespace cudf diff --git a/cpp/src/reductions/var.cu b/cpp/src/reductions/var.cu index 4d86918d6c6..d559531dc59 100644 --- a/cpp/src/reductions/var.cu +++ b/cpp/src/reductions/var.cu @@ -14,14 +14,16 @@ * limitations under the License. */ -#include +#include "compound.cuh" + #include -#include +#include #include namespace cudf { namespace reduction { +namespace detail { std::unique_ptr variance(column_view const& col, cudf::data_type const output_dtype, @@ -31,7 +33,7 @@ std::unique_ptr variance(column_view const& col, { // TODO: add cuda version check when the fix is available #if !defined(__CUDACC_DEBUG__) - using reducer = compound::detail::element_type_dispatcher; + using reducer = compound::detail::element_type_dispatcher; 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); @@ -42,5 +44,6 @@ std::unique_ptr variance(column_view const& col, #endif } +} // namespace detail } // namespace reduction } // namespace cudf