diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index a8e45b70572..e65ca2895c4 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -532,7 +532,6 @@ add_library( src/stream_compaction/apply_boolean_mask.cu src/stream_compaction/distinct.cu src/stream_compaction/distinct_count.cu - src/stream_compaction/distinct_reduce.cu src/stream_compaction/drop_nans.cu src/stream_compaction/drop_nulls.cu src/stream_compaction/stable_distinct.cu diff --git a/cpp/src/stream_compaction/distinct.cu b/cpp/src/stream_compaction/distinct.cu index b551df96765..8a7f6daa193 100644 --- a/cpp/src/stream_compaction/distinct.cu +++ b/cpp/src/stream_compaction/distinct.cu @@ -14,11 +14,11 @@ * limitations under the License. */ -#include "distinct_reduce.hpp" #include "stream_compaction_common.cuh" #include #include +#include #include #include #include @@ -39,6 +39,80 @@ namespace cudf { namespace detail { +namespace { +/** + * @brief Return the reduction identity used to initialize results of `hash_reduce_by_row`. + * + * @param keep A value of `duplicate_keep_option` type, must not be `KEEP_ANY`. + * @return The initial reduction value. + */ +auto constexpr reduction_init_value(duplicate_keep_option keep) +{ + switch (keep) { + case duplicate_keep_option::KEEP_FIRST: return std::numeric_limits::max(); + case duplicate_keep_option::KEEP_LAST: return std::numeric_limits::min(); + case duplicate_keep_option::KEEP_NONE: return size_type{0}; + default: CUDF_UNREACHABLE("This function should not be called with KEEP_ANY"); + } +} + +/** + * @brief The functor to find the first/last/all duplicate row for rows that compared equal. + */ +template +struct distinct_reduce_fn : reduce_by_row_fn_base { + duplicate_keep_option const keep; + + distinct_reduce_fn(MapView const& d_map, + KeyHasher const& d_hasher, + KeyEqual const& d_equal, + duplicate_keep_option const keep, + size_type* const d_output) + : reduce_by_row_fn_base{d_map, + d_hasher, + d_equal, + d_output}, + keep{keep} + { + } + + __device__ void operator()(size_type const idx) const + { + auto const out_ptr = this->get_output_ptr(idx); + + if (keep == duplicate_keep_option::KEEP_FIRST) { + // Store the smallest index of all rows that are equal. + atomicMin(out_ptr, idx); + } else if (keep == duplicate_keep_option::KEEP_LAST) { + // Store the greatest index of all rows that are equal. + atomicMax(out_ptr, idx); + } else { + // Count the number of rows in each group of rows that are compared equal. + atomicAdd(out_ptr, size_type{1}); + } + } +}; + +/** + * @brief The builder to construct an instance of `distinct_reduce_fn` functor base on the given + * value of the `duplicate_keep_option` member variable. + */ +struct reduce_func_builder { + duplicate_keep_option const keep; + + template + auto build(MapView const& d_map, + KeyHasher const& d_hasher, + KeyEqual const& d_equal, + size_type* const d_output) + { + return distinct_reduce_fn{ + d_map, d_hasher, d_equal, keep, d_output}; + } +}; + +} // namespace + rmm::device_uvector get_distinct_indices(table_view const& input, duplicate_keep_option keep, null_equality nulls_equal, @@ -97,16 +171,21 @@ rmm::device_uvector get_distinct_indices(table_view const& input, } // For other keep options, reduce by row on rows that compare equal. - auto const reduction_results = distinct_reduce(map, - std::move(preprocessed_input), - input.num_rows(), - has_nulls, - has_nested_columns, - keep, - nulls_equal, - nans_equal, - stream, - rmm::mr::get_current_device_resource()); + // Depending on the `keep` parameter, the reduction operation for each row group is: + // - If `keep == KEEP_FIRST`: min of row indices in the group. + // - If `keep == KEEP_LAST`: max of row indices in the group. + // - If `keep == KEEP_NONE`: count of equivalent rows (group size). + auto const reduction_results = hash_reduce_by_row(map, + std::move(preprocessed_input), + input.num_rows(), + has_nulls, + has_nested_columns, + nulls_equal, + nans_equal, + reduce_func_builder{keep}, + reduction_init_value(keep), + stream, + mr); // Extract the desired output indices from reduction results. auto const map_end = [&] { diff --git a/cpp/src/stream_compaction/distinct_reduce.cu b/cpp/src/stream_compaction/distinct_reduce.cu deleted file mode 100644 index 64d29ae2ff0..00000000000 --- a/cpp/src/stream_compaction/distinct_reduce.cu +++ /dev/null @@ -1,109 +0,0 @@ -/* - * 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. - * 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 "distinct_reduce.hpp" - -#include - -namespace cudf::detail { - -namespace { -/** - * @brief The functor to find the first/last/all duplicate row for rows that compared equal. - */ -template -struct distinct_reduce_fn : reduce_by_row_fn_base { - duplicate_keep_option const keep; - - distinct_reduce_fn(MapView const& d_map, - KeyHasher const& d_hasher, - KeyEqual const& d_equal, - duplicate_keep_option const keep, - size_type* const d_output) - : reduce_by_row_fn_base{d_map, - d_hasher, - d_equal, - d_output}, - keep{keep} - { - } - - __device__ void operator()(size_type const idx) const - { - auto const out_ptr = this->get_output_ptr(idx); - - if (keep == duplicate_keep_option::KEEP_FIRST) { - // Store the smallest index of all rows that are equal. - atomicMin(out_ptr, idx); - } else if (keep == duplicate_keep_option::KEEP_LAST) { - // Store the greatest index of all rows that are equal. - atomicMax(out_ptr, idx); - } else { - // Count the number of rows in each group of rows that are compared equal. - atomicAdd(out_ptr, size_type{1}); - } - } -}; - -/** - * @brief The builder to construct an instance of `distinct_reduce_fn` functor base on the given - * value of the `duplicate_keep_option` member variable. - */ -struct reduce_func_builder { - duplicate_keep_option const keep; - - template - auto build(MapView const& d_map, - KeyHasher const& d_hasher, - KeyEqual const& d_equal, - size_type* const d_output) - { - return distinct_reduce_fn{ - d_map, d_hasher, d_equal, keep, d_output}; - } -}; - -} // namespace - -rmm::device_uvector distinct_reduce( - hash_map_type const& map, - std::shared_ptr const preprocessed_input, - size_type num_rows, - cudf::nullate::DYNAMIC has_nulls, - bool has_nested_columns, - duplicate_keep_option keep, - null_equality nulls_equal, - nan_equality nans_equal, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) -{ - CUDF_EXPECTS(keep != duplicate_keep_option::KEEP_ANY, - "This function should not be called with KEEP_ANY"); - - return hash_reduce_by_row(map, - preprocessed_input, - num_rows, - has_nulls, - has_nested_columns, - nulls_equal, - nans_equal, - reduce_func_builder{keep}, - reduction_init_value(keep), - stream, - mr); -} - -} // namespace cudf::detail diff --git a/cpp/src/stream_compaction/distinct_reduce.hpp b/cpp/src/stream_compaction/distinct_reduce.hpp deleted file mode 100644 index 236b6c860c3..00000000000 --- a/cpp/src/stream_compaction/distinct_reduce.hpp +++ /dev/null @@ -1,87 +0,0 @@ -/* - * 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. - * 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 "stream_compaction_common.hpp" - -#include -#include -#include - -#include -#include - -namespace cudf::detail { - -/** - * @brief Return the reduction identity used to initialize results of `hash_reduce_by_row`. - * - * @param keep A value of `duplicate_keep_option` type, must not be `KEEP_ANY`. - * @return The initial reduction value. - */ -auto constexpr reduction_init_value(duplicate_keep_option keep) -{ - switch (keep) { - case duplicate_keep_option::KEEP_FIRST: return std::numeric_limits::max(); - case duplicate_keep_option::KEEP_LAST: return std::numeric_limits::min(); - case duplicate_keep_option::KEEP_NONE: return size_type{0}; - default: CUDF_UNREACHABLE("This function should not be called with KEEP_ANY"); - } -} - -/** - * @brief Perform a reduction on groups of rows that are compared equal. - * - * This is essentially a reduce-by-key operation with keys are non-contiguous rows and are compared - * equal. A hash table is used to find groups of equal rows. - * - * Depending on the `keep` parameter, the reduction operation for each row group is: - * - If `keep == KEEP_FIRST`: min of row indices in the group. - * - If `keep == KEEP_LAST`: max of row indices in the group. - * - If `keep == KEEP_NONE`: count of equivalent rows (group size). - * - * Note that this function is not needed when `keep == KEEP_NONE`. - * - * At the beginning of the operation, the entire output array is filled with a value given by - * the `reduction_init_value()` function. Then, the reduction result for each row group is written - * into the output array at the index of an unspecified row in the group. - * - * @param map The auxiliary map to perform reduction - * @param preprocessed_input The preprocessed of the input rows for computing row hashing and row - * comparisons - * @param num_rows The number of all input rows - * @param has_nulls Indicate whether the input rows has any nulls at any nested levels - * @param has_nested_columns Indicates whether the input table has any nested columns - * @param keep The parameter to determine what type of reduction to perform - * @param nulls_equal Flag to specify whether null elements should be considered as equal - * @param nans_equal Flag to specify whether NaN values in floating point column should be - * considered equal. - * @param stream CUDA stream used for device memory operations and kernel launches - * @param mr Device memory resource used to allocate the returned vector - * @return A device_uvector containing the reduction results - */ -rmm::device_uvector distinct_reduce( - hash_map_type const& map, - std::shared_ptr const preprocessed_input, - size_type num_rows, - cudf::nullate::DYNAMIC has_nulls, - bool has_nested_columns, - duplicate_keep_option keep, - null_equality nulls_equal, - nan_equality nans_equal, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr); - -} // namespace cudf::detail