From c13642623b01ec0f7ca223a3afcb7a6e33560321 Mon Sep 17 00:00:00 2001 From: Lawrence Mitchell Date: Mon, 26 Feb 2024 15:41:48 +0000 Subject: [PATCH] Move struct to implementation header Unify template parameter for stable sort as a sort_method enum member rather than bool. --- cpp/include/cudf/detail/sorting.hpp | 33 ---------------------------- cpp/src/sort/segmented_sort_impl.cuh | 9 +++----- cpp/src/sort/sort.cu | 15 ++++++------- cpp/src/sort/sort_column.cu | 15 +++++++------ cpp/src/sort/sort_column_impl.cuh | 14 +++++++----- cpp/src/sort/sort_impl.cuh | 7 +++--- cpp/src/sort/stable_sort.cu | 16 ++++++-------- cpp/src/sort/stable_sort_column.cu | 15 +++++++------ 8 files changed, 45 insertions(+), 79 deletions(-) diff --git a/cpp/include/cudf/detail/sorting.hpp b/cpp/include/cudf/detail/sorting.hpp index 92bc4117a27..97cc054da57 100644 --- a/cpp/include/cudf/detail/sorting.hpp +++ b/cpp/include/cudf/detail/sorting.hpp @@ -19,14 +19,8 @@ #include #include #include -#include -#include #include -#include - -#include -#include #include #include @@ -34,33 +28,6 @@ namespace cudf { namespace detail { -template -struct inplace_column_sort_fn { - template ()>* = nullptr> - void operator()(mutable_column_view& col, bool ascending, rmm::cuda_stream_view stream) const - { - CUDF_EXPECTS(!col.has_nulls(), "Nulls not supported for in-place sort"); - auto const do_sort = [&](auto const cmp) { - if constexpr (stable) { - thrust::stable_sort(rmm::exec_policy(stream), col.begin(), col.end(), cmp); - } else { - thrust::sort(rmm::exec_policy(stream), col.begin(), col.end(), cmp); - } - }; - if (ascending) { - do_sort(thrust::less()); - } else { - do_sort(thrust::greater()); - } - } - - template ()>* = nullptr> - void operator()(mutable_column_view&, bool, rmm::cuda_stream_view) const - { - CUDF_FAIL("Column type must be relationally comparable and fixed-width"); - } -}; - /** * @copydoc cudf::sorted_order * diff --git a/cpp/src/sort/segmented_sort_impl.cuh b/cpp/src/sort/segmented_sort_impl.cuh index 5d11bf055f1..0c927bbf963 100644 --- a/cpp/src/sort/segmented_sort_impl.cuh +++ b/cpp/src/sort/segmented_sort_impl.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. + * Copyright (c) 2021-2024, 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,6 +14,8 @@ * limitations under the License. */ +#include "common_sort_impl.cuh" + #include #include #include @@ -29,11 +31,6 @@ namespace cudf { namespace detail { -/** - * @brief The enum specifying which sorting method to use (stable or unstable). - */ -enum class sort_method { STABLE, UNSTABLE }; - /** * @brief Functor performs faster segmented sort on eligible columns */ diff --git a/cpp/src/sort/sort.cu b/cpp/src/sort/sort.cu index d185eb3c023..7fbef98f8b8 100644 --- a/cpp/src/sort/sort.cu +++ b/cpp/src/sort/sort.cu @@ -14,6 +14,7 @@ * limitations under the License. */ +#include "common_sort_impl.cuh" #include "sort_impl.cuh" #include @@ -37,7 +38,7 @@ std::unique_ptr sorted_order(table_view const& input, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - return sorted_order(input, column_order, null_precedence, stream, mr); + return sorted_order(input, column_order, null_precedence, stream, mr); } std::unique_ptr sort_by_key(table_view const& values, @@ -68,14 +69,12 @@ std::unique_ptr
sort(table_view const& input, rmm::mr::device_memory_resource* mr) { // fast-path sort conditions: single, non-floating-point, fixed-width column with no nulls - if (input.num_columns() == 1 && !input.column(0).has_nulls() && - cudf::is_fixed_width(input.column(0).type()) && - !cudf::is_floating_point(input.column(0).type())) { - auto output = std::make_unique(input.column(0), stream, mr); - auto view = output->mutable_view(); - bool ascending = (column_order.empty() ? true : column_order.front() == order::ASCENDING); + if (can_use_fast_sort(input)) { + auto output = std::make_unique(input.column(0), stream, mr); + auto view = output->mutable_view(); + auto order = (column_order.empty() ? order::ASCENDING : column_order.front()); cudf::type_dispatcher( - output->type(), inplace_column_sort_fn{}, view, ascending, stream); + output->type(), inplace_column_sort_fn{}, view, order, stream); std::vector> columns; columns.emplace_back(std::move(output)); return std::make_unique
(std::move(columns)); diff --git a/cpp/src/sort/sort_column.cu b/cpp/src/sort/sort_column.cu index 9df04251e93..7db44476988 100644 --- a/cpp/src/sort/sort_column.cu +++ b/cpp/src/sort/sort_column.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. + * Copyright (c) 2021-2024, 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,6 +14,7 @@ * limitations under the License. */ +#include "common_sort_impl.cuh" #include "sort_column_impl.cuh" #include @@ -30,11 +31,11 @@ namespace detail { * sorted_order(column_view&,order,null_order,rmm::cuda_stream_view,rmm::mr::device_memory_resource*) */ template <> -std::unique_ptr sorted_order(column_view const& input, - order column_order, - null_order null_precedence, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +std::unique_ptr sorted_order(column_view const& input, + order column_order, + null_order null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { auto sorted_indices = cudf::make_numeric_column( data_type(type_to_id()), input.size(), mask_state::UNALLOCATED, stream, mr); @@ -42,7 +43,7 @@ std::unique_ptr sorted_order(column_view const& input, thrust::sequence( rmm::exec_policy(stream), indices_view.begin(), indices_view.end(), 0); cudf::type_dispatcher(input.type(), - column_sorted_order_fn{}, + column_sorted_order_fn{}, input, indices_view, column_order == order::ASCENDING, diff --git a/cpp/src/sort/sort_column_impl.cuh b/cpp/src/sort/sort_column_impl.cuh index 5abc6bdfadf..7af24f22b67 100644 --- a/cpp/src/sort/sort_column_impl.cuh +++ b/cpp/src/sort/sort_column_impl.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. + * Copyright (c) 2021-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -16,6 +16,8 @@ #pragma once +#include "common_sort_impl.cuh" + #include #include #include @@ -36,7 +38,7 @@ namespace detail { * This API offers fast sorting for primitive types. It cannot handle nested types and will not * consider `NaN` as equivalent to other `NaN`. * - * @tparam stable Whether to use stable sort + * @tparam method Whether to use stable sort * @param input Column to sort. The column data is not modified. * @param column_order Ascending or descending sort order * @param null_precedence How null rows are to be ordered @@ -45,7 +47,7 @@ namespace detail { * @param mr Device memory resource used to allocate the returned column's device memory * @return Sorted indices for the input column. */ -template +template std::unique_ptr sorted_order(column_view const& input, order column_order, null_order null_precedence, @@ -78,7 +80,7 @@ struct simple_comparator { null_order null_precedence{}; }; -template +template struct column_sorted_order_fn { /** * @brief Compile time check for allowing faster sort. @@ -121,7 +123,7 @@ struct column_sorted_order_fn { auto const do_sort = [&](auto const comp) { // Compiling `thrust::*sort*` APIs is expensive. // Thus, we should optimize that by using constexpr condition to only compile what we need. - if constexpr (stable) { + if constexpr (method == sort_method::STABLE) { thrust::stable_sort_by_key(rmm::exec_policy(stream), d_col.begin(), d_col.end(), @@ -165,7 +167,7 @@ struct column_sorted_order_fn { auto comp = simple_comparator{*keys, input.has_nulls(), ascending, null_precedence}; // Compiling `thrust::*sort*` APIs is expensive. // Thus, we should optimize that by using constexpr condition to only compile what we need. - if constexpr (stable) { + if constexpr (method == sort_method::STABLE) { thrust::stable_sort( rmm::exec_policy(stream), indices.begin(), indices.end(), comp); } else { diff --git a/cpp/src/sort/sort_impl.cuh b/cpp/src/sort/sort_impl.cuh index dcab49475c7..e0331d65053 100644 --- a/cpp/src/sort/sort_impl.cuh +++ b/cpp/src/sort/sort_impl.cuh @@ -16,6 +16,7 @@ #pragma once +#include "common_sort_impl.cuh" #include "sort_column_impl.cuh" #include @@ -30,7 +31,7 @@ namespace detail { * @tparam stable Whether to use stable sort * @param stream CUDA stream used for device memory operations and kernel launches */ -template +template std::unique_ptr sorted_order(table_view input, std::vector const& column_order, std::vector const& null_precedence, @@ -57,7 +58,7 @@ std::unique_ptr sorted_order(table_view input, auto const single_col = input.column(0); auto const col_order = column_order.empty() ? order::ASCENDING : column_order.front(); auto const null_prec = null_precedence.empty() ? null_order::BEFORE : null_precedence.front(); - return sorted_order(single_col, col_order, null_prec, stream, mr); + return sorted_order(single_col, col_order, null_prec, stream, mr); } std::unique_ptr sorted_indices = cudf::make_numeric_column( @@ -71,7 +72,7 @@ std::unique_ptr sorted_order(table_view input, auto const do_sort = [&](auto const comparator) { // Compiling `thrust::*sort*` APIs is expensive. // Thus, we should optimize that by using constexpr condition to only compile what we need. - if constexpr (stable) { + if constexpr (method == sort_method::STABLE) { thrust::stable_sort(rmm::exec_policy(stream), mutable_indices_view.begin(), mutable_indices_view.end(), diff --git a/cpp/src/sort/stable_sort.cu b/cpp/src/sort/stable_sort.cu index 2d17ca3eeb5..dcc6a07457e 100644 --- a/cpp/src/sort/stable_sort.cu +++ b/cpp/src/sort/stable_sort.cu @@ -14,6 +14,7 @@ * limitations under the License. */ +#include "common_sort_impl.cuh" #include "sort_impl.cuh" #include @@ -34,7 +35,7 @@ std::unique_ptr stable_sorted_order(table_view const& input, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - return sorted_order(input, column_order, null_precedence, stream, mr); + return sorted_order(input, column_order, null_precedence, stream, mr); } std::unique_ptr
stable_sort(table_view const& input, @@ -43,15 +44,12 @@ std::unique_ptr
stable_sort(table_view const& input, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - // fast-path sort conditions: single, non-floating-point, fixed-width column with no nulls - if (input.num_columns() == 1 && !input.column(0).has_nulls() && - cudf::is_fixed_width(input.column(0).type()) && - !cudf::is_floating_point(input.column(0).type())) { - auto output = std::make_unique(input.column(0), stream, mr); - auto view = output->mutable_view(); - bool ascending = (column_order.empty() ? true : column_order.front() == order::ASCENDING); + if (can_use_fast_sort(input)) { + auto output = std::make_unique(input.column(0), stream, mr); + auto view = output->mutable_view(); + auto order = (column_order.empty() ? order::ASCENDING : column_order.front()); cudf::type_dispatcher( - output->type(), inplace_column_sort_fn{}, view, ascending, stream); + output->type(), inplace_column_sort_fn{}, view, order, stream); std::vector> columns; columns.emplace_back(std::move(output)); return std::make_unique
(std::move(columns)); diff --git a/cpp/src/sort/stable_sort_column.cu b/cpp/src/sort/stable_sort_column.cu index be519ead951..25a6c92034a 100644 --- a/cpp/src/sort/stable_sort_column.cu +++ b/cpp/src/sort/stable_sort_column.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. + * Copyright (c) 2021-2024, 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,6 +14,7 @@ * limitations under the License. */ +#include "common_sort_impl.cuh" #include "sort_column_impl.cuh" #include @@ -30,11 +31,11 @@ namespace detail { * sorted_order(column_view&,order,null_order,rmm::cuda_stream_view,rmm::mr::device_memory_resource*) */ template <> -std::unique_ptr sorted_order(column_view const& input, - order column_order, - null_order null_precedence, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +std::unique_ptr sorted_order(column_view const& input, + order column_order, + null_order null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { auto sorted_indices = cudf::make_numeric_column( data_type(type_to_id()), input.size(), mask_state::UNALLOCATED, stream, mr); @@ -42,7 +43,7 @@ std::unique_ptr sorted_order(column_view const& input, thrust::sequence( rmm::exec_policy(stream), indices_view.begin(), indices_view.end(), 0); cudf::type_dispatcher(input.type(), - column_sorted_order_fn{}, + column_sorted_order_fn{}, input, indices_view, column_order == order::ASCENDING,