diff --git a/cpp/include/cudf/detail/sorting.hpp b/cpp/include/cudf/detail/sorting.hpp index 8f92b66d5fa..97cc054da57 100644 --- a/cpp/include/cudf/detail/sorting.hpp +++ b/cpp/include/cudf/detail/sorting.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2023, NVIDIA CORPORATION. + * Copyright (c) 2019-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. @@ -150,5 +150,16 @@ std::unique_ptr sort(table_view const& values, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr); +/** + * @copydoc cudf::stable_sort + * + * @param stream CUDA stream used for device memory operations and kernel launches. + */ +std::unique_ptr
stable_sort(table_view const& values, + std::vector const& column_order, + std::vector const& null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); + } // namespace detail } // namespace cudf diff --git a/cpp/include/cudf/sorting.hpp b/cpp/include/cudf/sorting.hpp index e4e803b2d3c..42bcb5da8e3 100644 --- a/cpp/include/cudf/sorting.hpp +++ b/cpp/include/cudf/sorting.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2023, NVIDIA CORPORATION. + * Copyright (c) 2019-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. @@ -115,6 +115,18 @@ std::unique_ptr
sort( rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +/** + * @brief Performs a stable lexicographic sort of the rows of a table + * + * @copydoc cudf::sort + */ +std::unique_ptr
stable_sort( + table_view const& input, + std::vector const& column_order = {}, + std::vector const& null_precedence = {}, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + /** * @brief Performs a key-value sort. * @@ -148,26 +160,7 @@ std::unique_ptr
sort_by_key( /** * @brief Performs a key-value stable sort. * - * Creates a new table that reorders the rows of `values` according to the - * lexicographic ordering of the rows of `keys`. - * - * The order of equivalent elements is guaranteed to be preserved. - * - * @throws cudf::logic_error if `values.num_rows() != keys.num_rows()`. - * - * @param values The table to reorder - * @param keys The table that determines the ordering - * @param column_order The desired order for each column in `keys`. Size must be - * equal to `keys.num_columns()` or empty. If empty, all columns are sorted in - * ascending order. - * @param null_precedence The desired order of a null element compared to other - * elements for each column in `keys`. Size must be equal to - * `keys.num_columns()` or empty. If empty, all columns will be sorted with - * `null_order::BEFORE`. - * @param stream CUDA stream used for device memory operations and kernel launches - * @param mr Device memory resource used to allocate the returned table's device memory - * @return The reordering of `values` determined by the lexicographic order of - * the rows of `keys`. + * @copydoc cudf::sort_by_key */ std::unique_ptr
stable_sort_by_key( table_view const& values, diff --git a/cpp/src/sort/common_sort_impl.cuh b/cpp/src/sort/common_sort_impl.cuh new file mode 100644 index 00000000000..745e2717304 --- /dev/null +++ b/cpp/src/sort/common_sort_impl.cuh @@ -0,0 +1,101 @@ +/* + * Copyright (c) 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. + * 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. + */ +#pragma once + +#include +#include +#include +#include + +#include +#include + +#include + +namespace cudf { +namespace detail { + +/** + * @brief The enum specifying which sorting method to use (stable or unstable). + */ +enum class sort_method : bool { STABLE, UNSTABLE }; + +/** + * @brief Functor performs a fast-path, in-place sort on eligible columns + * + * @tparam method Whether to use a stable or unstable sort. + */ +template +struct inplace_column_sort_fn { + /** + * @brief Check if fast-path, in-place sort is available for the given column + * + * @param column to check + * @return true if fast-path sort is available, false otherwise. + */ + static bool is_usable(column_view const& column) + { + return !column.has_nulls() && cudf::is_fixed_width(column.type()) && + !cudf::is_floating_point(column.type()); + } + /** + * @brief Check if fast-path, in-place sort is available for the given table + * + * @param table to check + * @return true if fast-path sort is available, false otherwise. + */ + static bool is_usable(table_view const& table) + { + return table.num_columns() == 1 && is_usable(table.column(0)); + } + + /** + * @brief Fast-path sort a column in place + * + * Precondition, is_usable(column) returned true + * + * @tparam T column data type. + * @param col Column to sort, modified in place. + * @param order Ascending or descending sort order. + * @param stream CUDA stream used for device memory operations and kernel launches + * + */ + template ()>* = nullptr> + void operator()(mutable_column_view& col, order order, rmm::cuda_stream_view stream) const + { + auto const do_sort = [&](auto const cmp) { + if constexpr (method == sort_method::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 (order == order::ASCENDING) { + do_sort(thrust::less()); + } else { + do_sort(thrust::greater()); + } + } + + template ()>* = nullptr> + void operator()(mutable_column_view&, order, rmm::cuda_stream_view) const + { + CUDF_FAIL("Column type must be relationally comparable and fixed-width"); + } +}; + +} // namespace detail +} // namespace cudf diff --git a/cpp/src/sort/segmented_sort_impl.cuh b/cpp/src/sort/segmented_sort_impl.cuh index 5d11bf055f1..796e178fecd 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,10 @@ * limitations under the License. */ +#pragma once + +#include "common_sort_impl.cuh" + #include #include #include @@ -29,11 +33,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 46edae798d4..adffc06ab93 100644 --- a/cpp/src/sort/sort.cu +++ b/cpp/src/sort/sort.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2023, NVIDIA CORPORATION. + * Copyright (c) 2019-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_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, @@ -61,47 +62,24 @@ std::unique_ptr
sort_by_key(table_view const& values, mr); } -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"); - if (ascending) { - thrust::sort(rmm::exec_policy(stream), col.begin(), col.end(), thrust::less()); - } else { - thrust::sort(rmm::exec_policy(stream), col.begin(), col.end(), 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"); - } -}; - std::unique_ptr
sort(table_view const& input, std::vector const& column_order, std::vector const& null_precedence, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - CUDF_FUNC_RANGE(); // 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 (inplace_column_sort_fn::is_usable(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)); } - return detail::sort_by_key( - input, input, column_order, null_precedence, cudf::get_default_stream(), mr); + return detail::sort_by_key(input, input, column_order, null_precedence, stream, mr); } } // namespace detail 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 5fae8db1a70..e0331d65053 100644 --- a/cpp/src/sort/sort_impl.cuh +++ b/cpp/src/sort/sort_impl.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. + * Copyright (c) 2020-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,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, @@ -39,7 +40,7 @@ std::unique_ptr sorted_order(table_view input, { if (input.num_rows() == 0 or input.num_columns() == 0) { return cudf::make_numeric_column( - data_type(type_to_id()), 0, mask_state::UNALLOCATED, stream); + data_type(type_to_id()), 0, mask_state::UNALLOCATED, stream, mr); } if (not column_order.empty()) { @@ -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 cf602dcf1a9..0bfe2cfef16 100644 --- a/cpp/src/sort/stable_sort.cu +++ b/cpp/src/sort/stable_sort.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2023, NVIDIA CORPORATION. + * Copyright (c) 2019-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_impl.cuh" #include @@ -34,7 +35,26 @@ 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, + std::vector const& column_order, + std::vector const& null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + if (inplace_column_sort_fn::is_usable(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, order, stream); + std::vector> columns; + columns.emplace_back(std::move(output)); + return std::make_unique
(std::move(columns)); + } + return detail::stable_sort_by_key(input, input, column_order, null_precedence, stream, mr); } std::unique_ptr
stable_sort_by_key(table_view const& values, @@ -69,6 +89,16 @@ std::unique_ptr stable_sorted_order(table_view const& input, return detail::stable_sorted_order(input, column_order, null_precedence, stream, mr); } +std::unique_ptr
stable_sort(table_view const& input, + std::vector const& column_order, + std::vector const& null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::stable_sort(input, column_order, null_precedence, stream, mr); +} + std::unique_ptr
stable_sort_by_key(table_view const& values, table_view const& keys, std::vector const& column_order, 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, diff --git a/cpp/tests/sort/stable_sort_tests.cpp b/cpp/tests/sort/stable_sort_tests.cpp index 71520ef007b..341f8317004 100644 --- a/cpp/tests/sort/stable_sort_tests.cpp +++ b/cpp/tests/sort/stable_sort_tests.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -34,12 +34,14 @@ void run_stable_sort_test(cudf::table_view input, cudf::column_view expected_sorted_indices, std::vector column_order = {}, - std::vector null_precedence = {}) + std::vector null_precedence = {}, + bool by_key = true) { - auto got_sort_by_key_table = cudf::sort_by_key(input, input, column_order, null_precedence); - auto expected_sort_by_key_table = cudf::gather(input, expected_sorted_indices); + auto got = by_key ? cudf::stable_sort_by_key(input, input, column_order, null_precedence) + : cudf::stable_sort(input, column_order, null_precedence); + auto expected = cudf::gather(input, expected_sorted_indices); - CUDF_TEST_EXPECT_TABLES_EQUAL(expected_sort_by_key_table->view(), got_sort_by_key_table->view()); + CUDF_TEST_EXPECT_TABLES_EQUAL(expected->view(), got->view()); } using TestTypes = cudf::test::Concat col3{{10, 40, 70, 10, 2, 10}, {1, 1, 0, 1, 1, 1}}; cudf::table_view input{{col1, col2, col3}}; - cudf::test::fixed_width_column_wrapper expected{{1, 0, 3, 5, 4, 2}}; std::vector column_order{ cudf::order::ASCENDING, cudf::order::ASCENDING, cudf::order::DESCENDING}; std::vector null_precedence{ cudf::null_order::AFTER, cudf::null_order::AFTER, cudf::null_order::AFTER}; + auto expected = std::is_same_v + // All the bools are true, and therefore don't affect sort order, + // so this is just the sort order of the nullable string column + ? cudf::test::fixed_width_column_wrapper{{0, 3, 5, 1, 4, 2}} + : cudf::test::fixed_width_column_wrapper{{1, 0, 3, 5, 4, 2}}; auto got = cudf::stable_sorted_order(input, column_order, null_precedence); - if (not std::is_same_v) { - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - - run_stable_sort_test(input, expected, column_order, null_precedence); - } else { - // for bools only validate that the null element landed at the back, since - // the rest of the values are equivalent and yields random sorted order. - auto to_host = [](cudf::column_view const& col) { - thrust::host_vector h_data(col.size()); - CUDF_CUDA_TRY(cudaMemcpy( - h_data.data(), col.data(), h_data.size() * sizeof(int32_t), cudaMemcpyDefault)); - return h_data; - }; - thrust::host_vector h_exp = to_host(expected); - thrust::host_vector h_got = to_host(got->view()); - EXPECT_EQ(h_exp[h_exp.size() - 1], h_got[h_got.size() - 1]); - - cudf::test::fixed_width_column_wrapper expected_for_bool{{0, 3, 5, 1, 4, 2}}; - run_stable_sort_test(input, expected_for_bool, column_order, null_precedence); - } + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); + + run_stable_sort_test(input, expected, column_order, null_precedence, false); + run_stable_sort_test(input, expected, column_order, null_precedence, true); +} + +TYPED_TEST(StableSort, SingleColumnNoNull) +{ + // This test exercises the "fast-path" single column sort. + using T = TypeParam; + // 0 1 2 3 4 5 6 7 8 9 + cudf::test::fixed_width_column_wrapper col{{7, 1, -2, 5, 1, 0, 1, -2, 0, 5}}; + cudf::table_view input{{col}}; + std::vector column_order{cudf::order::ASCENDING}; + auto expected = + std::is_same_v + ? cudf::test::fixed_width_column_wrapper{{8, 5, 0, 1, 2, 3, 4, 6, 7, 9}} + : std::is_unsigned_v + ? cudf::test::fixed_width_column_wrapper{{5, 8, 1, 4, 6, 3, 9, 0, 2, 7}} + : cudf::test::fixed_width_column_wrapper{{2, 7, 5, 8, 1, 4, 6, 3, 9, 0}}; + run_stable_sort_test(input, expected, column_order, {}, false); + run_stable_sort_test(input, expected, column_order, {}, true); +} + +TYPED_TEST(StableSort, SingleColumnWithNull) +{ + using T = TypeParam; + // 0 1 2 3 4 5 6 7 8 9 + cudf::test::fixed_width_column_wrapper col{{7, 1, -2, 5, 1, 0, 1, -2, 0, 5}, + {1, 1, 0, 0, 1, 0, 1, 0, 1, 0}}; + cudf::table_view input{{col}}; + std::vector column_order{cudf::order::ASCENDING}; + std::vector null_precedence{cudf::null_order::BEFORE}; + auto expected = + std::is_same_v + ? cudf::test::fixed_width_column_wrapper{{5, 2, 3, 7, 9, 8, 0, 1, 4, 6}} + : std::is_unsigned_v + ? cudf::test::fixed_width_column_wrapper{{5, 3, 9, 2, 7, 8, 1, 4, 6, 0}} + : cudf::test::fixed_width_column_wrapper{{2, 7, 5, 3, 9, 8, 1, 4, 6, 0}}; + run_stable_sort_test(input, expected, column_order, {}, false); + run_stable_sort_test(input, expected, column_order, {}, true); } TYPED_TEST(StableSort, WithNullMin) @@ -117,32 +144,19 @@ TYPED_TEST(StableSort, WithNullMin) cudf::test::fixed_width_column_wrapper col3{{10, 40, 70, 10, 2}, {1, 1, 0, 1, 1}}; cudf::table_view input{{col1, col2, col3}}; - cudf::test::fixed_width_column_wrapper expected{{2, 1, 0, 3, 4}}; std::vector column_order{ cudf::order::ASCENDING, cudf::order::ASCENDING, cudf::order::DESCENDING}; + auto expected = std::is_same_v + // All the bools are true, and therefore don't affect sort order, + // so this is just the sort order of the string column + ? cudf::test::fixed_width_column_wrapper{{2, 0, 3, 1, 4}} + : cudf::test::fixed_width_column_wrapper{{2, 1, 0, 3, 4}}; + auto got = cudf::stable_sorted_order(input, column_order); - auto got = cudf::stable_sorted_order(input, column_order); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - if (!std::is_same_v) { - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - - run_stable_sort_test(input, expected, column_order); - } else { - // for bools only validate that the null element landed at the front, since - // the rest of the values are equivalent and yields random sorted order. - auto to_host = [](cudf::column_view const& col) { - thrust::host_vector h_data(col.size()); - CUDF_CUDA_TRY(cudaMemcpy( - h_data.data(), col.data(), h_data.size() * sizeof(int32_t), cudaMemcpyDefault)); - return h_data; - }; - thrust::host_vector h_exp = to_host(expected); - thrust::host_vector h_got = to_host(got->view()); - EXPECT_EQ(h_exp.front(), h_got.front()); - - cudf::test::fixed_width_column_wrapper expected_for_bool{{2, 0, 3, 1, 4}}; - run_stable_sort_test(input, expected_for_bool, column_order); - } + run_stable_sort_test(input, expected, column_order, {}, false); + run_stable_sort_test(input, expected, column_order, {}, true); } TYPED_TEST(StableSort, WithAllValid) @@ -154,22 +168,19 @@ TYPED_TEST(StableSort, WithAllValid) cudf::test::fixed_width_column_wrapper col3{{10, 40, 70, 10, 2}}; cudf::table_view input{{col1, col2, col3}}; - cudf::test::fixed_width_column_wrapper expected{{2, 1, 0, 3, 4}}; std::vector column_order{ cudf::order::ASCENDING, cudf::order::ASCENDING, cudf::order::DESCENDING}; + auto expected = std::is_same_v + // All the bools are true, and therefore don't affect sort order, + // so this is just the sort order of the string column + ? cudf::test::fixed_width_column_wrapper{{2, 0, 3, 1, 4}} + : cudf::test::fixed_width_column_wrapper{{2, 1, 0, 3, 4}}; + auto got = cudf::stable_sorted_order(input, column_order); - auto got = cudf::stable_sorted_order(input, column_order); - - // Skip validating bools order. Valid true bools are all - // equivalent, and yield random order after thrust::sort - if (!std::is_same_v) { - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - run_stable_sort_test(input, expected, column_order); - } else { - cudf::test::fixed_width_column_wrapper expected_for_bool{{2, 0, 3, 1, 4}}; - run_stable_sort_test(input, expected_for_bool, column_order); - } + run_stable_sort_test(input, expected, column_order, {}, false); + run_stable_sort_test(input, expected, column_order, {}, true); } TYPED_TEST(StableSort, MisMatchInColumnOrderSize)