diff --git a/cpp/include/cudf/detail/reduction_functions.hpp b/cpp/include/cudf/detail/reduction_functions.hpp index ccec4bf8a6c..3a6113e66ce 100644 --- a/cpp/include/cudf/detail/reduction_functions.hpp +++ b/cpp/include/cudf/detail/reduction_functions.hpp @@ -192,6 +192,8 @@ std::unique_ptr 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 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`. diff --git a/cpp/src/reductions/all.cu b/cpp/src/reductions/all.cu index 3a076c3b780..b43df279393 100644 --- a/cpp/src/reductions/all.cu +++ b/cpp/src/reductions/all.cu @@ -88,11 +88,12 @@ std::unique_ptr 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{}, - col, - stream, - mr); + return cudf::type_dispatcher( + col.type(), + simple::detail::bool_result_element_dispatcher{}, + col, + stream, + mr); } } // namespace reduction diff --git a/cpp/src/reductions/any.cu b/cpp/src/reductions/any.cu index 1eb080cfe20..bad7d581255 100644 --- a/cpp/src/reductions/any.cu +++ b/cpp/src/reductions/any.cu @@ -88,11 +88,12 @@ std::unique_ptr 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{}, - col, - stream, - mr); + return cudf::type_dispatcher( + col.type(), + simple::detail::bool_result_element_dispatcher{}, + col, + stream, + mr); } } // namespace reduction diff --git a/cpp/src/reductions/compound.cuh b/cpp/src/reductions/compound.cuh index c60c819f8e2..89a95f5138c 100644 --- a/cpp/src/reductions/compound.cuh +++ b/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 std::unique_ptr compound_reduction(column_view const& col, @@ -61,19 +60,19 @@ std::unique_ptr compound_reduction(column_view const& col, auto it = thrust::make_transform_iterator( dcol->pair_begin(), compound_op.template get_null_replacing_element_transformer()); - result = detail::reduce( + result = cudf::reduction::detail::reduce( it, col.size(), compound_op, valid_count, ddof, stream, mr); } else { auto it = thrust::make_transform_iterator( dcol->begin(), compound_op.template get_element_transformer()); - result = detail::reduce( + result = cudf::reduction::detail::reduce( it, col.size(), compound_op, valid_count, ddof, stream, mr); } } else { auto it = thrust::make_transform_iterator( cudf::dictionary::detail::make_dictionary_pair_iterator(*dcol, col.has_nulls()), compound_op.template get_null_replacing_element_transformer()); - result = detail::reduce( + result = cudf::reduction::detail::reduce( 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 diff --git a/cpp/src/reductions/max.cu b/cpp/src/reductions/max.cu index dd283d86d3b..4adf35414dd 100644 --- a/cpp/src/reductions/max.cu +++ b/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 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{}, - col, - stream, - mr); + return cudf::type_dispatcher( + dispatch_type, + simple::detail::same_element_type_dispatcher{}, + col, + stream, + mr); } } // namespace reduction diff --git a/cpp/src/reductions/mean.cu b/cpp/src/reductions/mean.cu index ca341090b9f..e4b5f754b9b 100644 --- a/cpp/src/reductions/mean.cu +++ b/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,7 +13,6 @@ * See the License for the specific language governing permissions and * limitations under the License. */ -// The translation unit for reduction `mean` #include #include @@ -21,14 +20,20 @@ #include -std::unique_ptr 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 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; + 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 is not used for mean*/ 1, stream, mr); } + +} // namespace reduction +} // namespace cudf diff --git a/cpp/src/reductions/min.cu b/cpp/src/reductions/min.cu index 5e1301b2904..ac9bdfe9cdc 100644 --- a/cpp/src/reductions/min.cu +++ b/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 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{}, - col, - stream, - mr); + return cudf::type_dispatcher( + dispatch_type, + simple::detail::same_element_type_dispatcher{}, + col, + stream, + mr); } } // namespace reduction diff --git a/cpp/src/reductions/product.cu b/cpp/src/reductions/product.cu index 30342bc4728..5caf498712a 100644 --- a/cpp/src/reductions/product.cu +++ b/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 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{}, + simple::detail::element_type_dispatcher{}, col, output_dtype, stream, diff --git a/cpp/src/reductions/simple.cuh b/cpp/src/reductions/simple.cuh index e5303246452..807462d742f 100644 --- a/cpp/src/reductions/simple.cuh +++ b/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 simple_reduction(column_view const& col, if (col.has_nulls()) { auto f = simple_op.template get_null_replacing_element_transformer(); auto it = thrust::make_transform_iterator(dcol->pair_begin(), 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(); auto it = thrust::make_transform_iterator(dcol->begin(), 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 fixed_point_reduction(column_view const& col, if (col.has_nulls()) { auto f = simple_op.template get_null_replacing_element_transformer(); auto it = thrust::make_transform_iterator(dcol->pair_begin(), 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(); auto it = thrust::make_transform_iterator(dcol->begin(), 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 dictionary_reduction(column_view const& col, auto p = cudf::dictionary::detail::make_dictionary_pair_iterator(*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(col, stream, mr); + return simple_reduction(col, stream, mr); } - auto index = simple::simple_reduction( + auto index = simple_reduction( 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 diff --git a/cpp/src/reductions/std.cu b/cpp/src/reductions/std.cu index 3c7a05abd4e..bb29e5cd030 100644 --- a/cpp/src/reductions/std.cu +++ b/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,7 +13,6 @@ * See the License for the specific language governing permissions and * limitations under the License. */ -// The translation unit for reduction `standard deviation` #include #include @@ -21,21 +20,19 @@ #include -// @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::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 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; + 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); @@ -45,3 +42,6 @@ std::unique_ptr cudf::reduction::standard_deviation( CUDF_FAIL("var/std reductions are not supported at debug build."); #endif } + +} // namespace reduction +} // namespace cudf diff --git a/cpp/src/reductions/sum.cu b/cpp/src/reductions/sum.cu index 8bc157668f4..2db19939bd5 100644 --- a/cpp/src/reductions/sum.cu +++ b/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 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{}, + simple::detail::element_type_dispatcher{}, col, output_dtype, stream, diff --git a/cpp/src/reductions/sum_of_squares.cu b/cpp/src/reductions/sum_of_squares.cu index eca6aa0d1d9..a3e9368bb02 100644 --- a/cpp/src/reductions/sum_of_squares.cu +++ b/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 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{}, + simple::detail::element_type_dispatcher{}, col, output_dtype, stream, diff --git a/cpp/src/reductions/var.cu b/cpp/src/reductions/var.cu index 2565e472661..2df653858b0 100644 --- a/cpp/src/reductions/var.cu +++ b/cpp/src/reductions/var.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. @@ -14,27 +14,24 @@ * limitations under the License. */ -// The translation unit for reduction `variance` - #include #include #include #include -// @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::reduction::variance(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 variance(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; + 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); @@ -44,3 +41,6 @@ std::unique_ptr cudf::reduction::variance(column_view const& col, CUDF_FAIL("var/std reductions are not supported at debug build."); #endif } + +} // namespace reduction +} // namespace cudf