diff --git a/cpp/benchmarks/sort/nested_types_common.hpp b/cpp/benchmarks/sort/nested_types_common.hpp index fabef3a7a51..e0626b1b96f 100644 --- a/cpp/benchmarks/sort/nested_types_common.hpp +++ b/cpp/benchmarks/sort/nested_types_common.hpp @@ -28,17 +28,23 @@ #include -inline std::unique_ptr create_lists_data(nvbench::state& state) +inline std::unique_ptr create_lists_data(nvbench::state& state, + cudf::size_type const num_columns = 1, + cudf::size_type const min_val = 0, + cudf::size_type const max_val = 5) { const size_t size_bytes(state.get_int64("size_bytes")); const cudf::size_type depth{static_cast(state.get_int64("depth"))}; auto const null_frequency{state.get_float64("null_frequency")}; data_profile table_profile; - table_profile.set_distribution_params(cudf::type_id::LIST, distribution_id::UNIFORM, 0, 5); + table_profile.set_distribution_params( + cudf::type_id::LIST, distribution_id::UNIFORM, min_val, max_val); table_profile.set_list_depth(depth); table_profile.set_null_probability(null_frequency); - return create_random_table({cudf::type_id::LIST}, table_size_bytes{size_bytes}, table_profile); + return create_random_table(std::vector(num_columns, cudf::type_id::LIST), + table_size_bytes{size_bytes}, + table_profile); } inline std::unique_ptr create_structs_data(nvbench::state& state, diff --git a/cpp/benchmarks/sort/sort_lists.cpp b/cpp/benchmarks/sort/sort_lists.cpp index b55b60f5ec9..3cab60a29ec 100644 --- a/cpp/benchmarks/sort/sort_lists.cpp +++ b/cpp/benchmarks/sort/sort_lists.cpp @@ -20,18 +20,82 @@ #include -void nvbench_sort_lists(nvbench::state& state) +namespace { +constexpr cudf::size_type min_val = 0; +constexpr cudf::size_type max_val = 100; + +void sort_multiple_lists(nvbench::state& state) +{ + auto const num_columns = static_cast(state.get_int64("num_columns")); + auto const input_table = create_lists_data(state, num_columns, min_val, max_val); + auto const stream = cudf::get_default_stream(); + + state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value())); + state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) { + cudf::detail::sorted_order( + *input_table, {}, {}, stream, rmm::mr::get_current_device_resource()); + }); +} + +void sort_lists_of_structs(nvbench::state& state) { - auto const table = create_lists_data(state); + auto const num_columns = static_cast(state.get_int64("num_columns")); + auto const lists_table = create_lists_data(state, num_columns, min_val, max_val); + // After having a table of (multiple) lists columns, convert those lists columns into lists of + // structs columns. The children of these structs columns are also children of the original lists + // columns. + // Such resulted lists-of-structs columns are very similar to the original lists-of-integers + // columns so their benchmarks can be somewhat comparable. + std::vector lists_of_structs; + for (auto const& col : lists_table->view()) { + auto const child = col.child(cudf::lists_column_view::child_column_index); + + // Put the child column under a struct column having the same null mask/null count. + auto const new_child = cudf::column_view{cudf::data_type{cudf::type_id::STRUCT}, + child.size(), + nullptr, + child.null_mask(), + child.null_count(), + child.offset(), + {child}}; + auto const converted_col = + cudf::column_view{cudf::data_type{cudf::type_id::LIST}, + col.size(), + nullptr, + col.null_mask(), + col.null_count(), + col.offset(), + {col.child(cudf::lists_column_view::offsets_column_index), new_child}}; + lists_of_structs.push_back(converted_col); + } + + auto const input_table = cudf::table_view{lists_of_structs}; + auto const stream = cudf::get_default_stream(); + + state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value())); state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) { rmm::cuda_stream_view stream_view{launch.get_stream()}; - cudf::detail::sorted_order(*table, {}, {}, stream_view, rmm::mr::get_current_device_resource()); + cudf::detail::sorted_order(input_table, {}, {}, stream, rmm::mr::get_current_device_resource()); }); } +} // namespace + +void nvbench_sort_lists(nvbench::state& state) +{ + const auto has_lists_of_structs = state.get_int64("lists_of_structs") > 0; + if (has_lists_of_structs) { + sort_lists_of_structs(state); + } else { + sort_multiple_lists(state); + } +} + NVBENCH_BENCH(nvbench_sort_lists) .set_name("sort_list") .add_int64_power_of_two_axis("size_bytes", {10, 18, 24, 28}) .add_int64_axis("depth", {1, 4}) + .add_int64_axis("num_columns", {1}) + .add_int64_axis("lists_of_structs", {0, 1}) .add_float64_axis("null_frequency", {0, 0.2}); diff --git a/cpp/include/cudf/detail/sorting.hpp b/cpp/include/cudf/detail/sorting.hpp index 9c936e6c659..8f92b66d5fa 100644 --- a/cpp/include/cudf/detail/sorting.hpp +++ b/cpp/include/cudf/detail/sorting.hpp @@ -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. @@ -16,6 +16,7 @@ #pragma once +#include #include #include @@ -61,6 +62,20 @@ std::unique_ptr sort_by_key(table_view const& values, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr); +/** + * @copydoc cudf::rank + * + * @param[in] stream CUDA stream used for device memory operations and kernel launches. + */ +std::unique_ptr rank(column_view const& input, + rank_method method, + order column_order, + null_policy null_handling, + null_order null_precedence, + bool percentage, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); + /** * @copydoc cudf::stable_sort_by_key * diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index a549e725605..3e37bd53972 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -679,16 +679,22 @@ struct preprocessed_table { * @brief Preprocess table for use with lexicographical comparison * * Sets up the table for use with lexicographical comparison. The resulting preprocessed table can - * be passed to the constructor of `lexicographic::self_comparator` to avoid preprocessing again. + * be passed to the constructor of `lexicographic::self_comparator` or + * `lexicographic::two_table_comparator` to avoid preprocessing again. + * + * Note that the output of this factory function should not be used in `two_table_comparator` if + * the input table contains lists-of-structs. In such cases, please use the overload + * `preprocessed_table::create(table_view const&, table_view const&,...)` to preprocess both input + * tables at the same time. * * @param table The table to preprocess * @param column_order Optional, host array the same length as a row that indicates the desired - * ascending/descending order of each column in a row. If empty, it is assumed all columns are - * sorted in ascending order. - * @param null_precedence Optional, device array the same length as a row and indicates how null - * values compare to all other for every column. If it is nullptr, then null precedence would be - * `null_order::BEFORE` for all columns. - * @param stream The stream to launch kernels and h->d copies on while preprocessing. + * ascending/descending order of each column in a row. If empty, it is assumed all columns + * are sorted in ascending order. + * @param null_precedence Optional, an array having the same length as the number of columns in + * the input tables that indicates how null values compare to all other. If it is empty, + * the order `null_order::BEFORE` will be used for all columns. + * @param stream The stream to launch kernels and h->d copies on while preprocessing * @return A shared pointer to a preprocessed table */ static std::shared_ptr create(table_view const& table, @@ -696,10 +702,66 @@ struct preprocessed_table { host_span null_precedence, rmm::cuda_stream_view stream); + /** + * @brief Preprocess tables for use with lexicographical comparison + * + * Sets up the tables for use with lexicographical comparison. The resulting preprocessed tables + * can be passed to the constructor of `lexicographic::self_comparator` or + * `lexicographic::two_table_comparator` to avoid preprocessing again. + * + * This factory function performs some extra operations to guarantee that its output can be used + * in `two_table_comparator` for all cases. + * + * @param lhs The lhs table to preprocess + * @param rhs The rhs table to preprocess + * @param column_order Optional, host array the same length as a row that indicates the desired + * ascending/descending order of each column in a row. If empty, it is assumed all columns + * are sorted in ascending order. + * @param null_precedence Optional, an array having the same length as the number of columns in + * the input tables that indicates how null values compare to all other. If it is empty, + * the order `null_order::BEFORE` will be used for all columns. + * @param stream The stream to launch kernels and h->d copies on while preprocessing + * @return A pair of shared pointers to the preprocessed tables + */ + static std::pair, std::shared_ptr> create( + table_view const& lhs, + table_view const& rhs, + host_span column_order, + host_span null_precedence, + rmm::cuda_stream_view stream); + private: friend class self_comparator; ///< Allow self_comparator to access private members friend class two_table_comparator; ///< Allow two_table_comparator to access private members + /** + * @brief Create the output preprocessed table from intermediate preprocessing results + * + * @param preprocessed_input The table resulted from preprocessing + * @param verticalized_col_depths The depths of each column resulting from decomposing struct + * columns in the original input table + * @param transformed_columns Store the intermediate columns generated from transforming + * nested children columns into integers columns using `cudf::rank()` + * @param column_order Optional, host array the same length as a row that indicates the desired + * ascending/descending order of each column in a row. If empty, it is assumed all columns + * are sorted in ascending order. + * @param null_precedence Optional, an array having the same length as the number of columns in + * the input tables that indicates how null values compare to all other. If it is empty, + * the order `null_order::BEFORE` will be used for all columns. + * @param has_ranked_children Flag indicating if the input table was preprocessed to transform + * any nested child column into an integer column using `cudf::rank` + * @param stream The stream to launch kernels and h->d copies on while preprocessing + * @return A shared pointer to a preprocessed table + */ + static std::shared_ptr create( + table_view const& preprocessed_input, + std::vector&& verticalized_col_depths, + std::vector>&& transformed_columns, + host_span column_order, + host_span null_precedence, + bool has_ranked_children, + rmm::cuda_stream_view stream); + /** * @brief Construct a preprocessed table for use with lexicographical comparison * @@ -708,43 +770,39 @@ struct preprocessed_table { * * @param table The table to preprocess * @param column_order Optional, device array the same length as a row that indicates the desired - * ascending/descending order of each column in a row. If empty, it is assumed all columns are - * sorted in ascending order. + * ascending/descending order of each column in a row. If empty, it is assumed all columns + * are sorted in ascending order. * @param null_precedence Optional, device array the same length as a row and indicates how null - * values compare to all other for every column. If it is nullptr, then null precedence would be - * `null_order::BEFORE` for all columns. + * values compare to all other for every column. If it is nullptr, then null precedence + * would be `null_order::BEFORE` for all columns. * @param depths The depths of each column resulting from decomposing struct columns. * @param dremel_data The dremel data for each list column. The length of this object is the - * number of list columns in the table. + * number of list columns in the table. * @param dremel_device_views Device views into the dremel_data structs contained in the - * `dremel_data` parameter. For columns that are not list columns, this uvector will should - * contain an empty `dremel_device_view`. As such, this uvector has as many elements as there are - * columns in the table (unlike the `dremel_data` parameter, which is only as long as the number - * of list columns). + * `dremel_data` parameter. For columns that are not list columns, this uvector will should + * contain an empty `dremel_device_view`. As such, this uvector has as many elements as + * there are columns in the table (unlike the `dremel_data` parameter, which is only as + * long as the number of list columns). + * @param transformed_columns Store the intermediate columns generated from transforming + * nested children columns into integers columns using `cudf::rank()` + * @param has_ranked_children Flag indicating if the input table was preprocessed to transform + * any lists-of-structs column having floating-point children using `cudf::rank` */ preprocessed_table(table_device_view_owner&& table, rmm::device_uvector&& column_order, rmm::device_uvector&& null_precedence, rmm::device_uvector&& depths, std::vector&& dremel_data, - rmm::device_uvector&& dremel_device_views) - : _t(std::move(table)), - _column_order(std::move(column_order)), - _null_precedence(std::move(null_precedence)), - _depths(std::move(depths)), - _dremel_data(std::move(dremel_data)), - _dremel_device_views(std::move(dremel_device_views)){}; + rmm::device_uvector&& dremel_device_views, + std::vector>&& transformed_columns, + bool has_ranked_children); preprocessed_table(table_device_view_owner&& table, rmm::device_uvector&& column_order, rmm::device_uvector&& null_precedence, - rmm::device_uvector&& depths) - : _t(std::move(table)), - _column_order(std::move(column_order)), - _null_precedence(std::move(null_precedence)), - _depths(std::move(depths)), - _dremel_data{}, - _dremel_device_views{} {}; + rmm::device_uvector&& depths, + std::vector>&& transformed_columns, + bool has_ranked_children); /** * @brief Implicit conversion operator to a `table_device_view` of the preprocessed table. @@ -800,6 +858,16 @@ struct preprocessed_table { } } + template + void check_physical_element_comparator() + { + if constexpr (!std::is_same_v) { + CUDF_EXPECTS(!_has_ranked_children, + "The input table has nested type children and they were transformed using a " + "different type of physical element comparator."); + } + } + private: table_device_view_owner const _t; rmm::device_uvector const _column_order; @@ -809,6 +877,14 @@ struct preprocessed_table { // Dremel encoding of list columns used for the comparison algorithm std::optional> _dremel_data; std::optional> _dremel_device_views; + + // Intermediate columns generated from transforming nested children columns into + // integers columns using `cudf::rank()`, need to be kept alive. + std::vector> _transformed_columns; + + // Flag to record if the input table was preprocessed to transform any nested children column(s) + // into integer column(s) using `cudf::rank`. + bool const _has_ranked_children; }; /** @@ -832,13 +908,13 @@ class self_comparator { * * @param t The table to compare * @param column_order Optional, host array the same length as a row that indicates the desired - * ascending/descending order of each column in a row. If empty, it is assumed all columns are - * sorted in ascending order. + * ascending/descending order of each column in a row. If empty, it is assumed all columns + * are sorted in ascending order. * @param null_precedence Optional, device array the same length as a row and indicates how null - * values compare to all other for every column. If empty, then null precedence would be - * `null_order::BEFORE` for all columns. + * values compare to all other for every column. If empty, then null precedence would be + * `null_order::BEFORE` for all columns. * @param stream The stream to construct this object on. Not the stream that will be used for - * comparisons using this object. + * comparisons using this object. */ self_comparator(table_view const& t, host_span column_order = {}, @@ -867,9 +943,9 @@ class self_comparator { * `F(i,j)` returns true if and only if row `i` compares lexicographically less than row `j`. * * @note The operator overloads in sub-class `element_comparator` are templated via the - * `type_dispatcher` to help select an overload instance for each column in a table. - * So, `cudf::is_nested` will return `true` if the table has nested-type columns, - * but it will be a runtime error if template parameter `has_nested_columns != true`. + * `type_dispatcher` to help select an overload instance for each column in a table. + * So, `cudf::is_nested` will return `true` if the table has nested-type columns, + * but it will be a runtime error if template parameter `has_nested_columns != true`. * * @tparam has_nested_columns compile-time optimization for primitive types. * This template parameter is to be used by the developer by querying @@ -878,8 +954,11 @@ class self_comparator { * overloads for primitive types. * @tparam Nullate A cudf::nullate type describing whether to check for nulls. * @tparam PhysicalElementComparator A relational comparator functor that compares individual - * values rather than logical elements, defaults to `NaN` aware relational comparator that - * evaluates `NaN` as greater than all other values. + * values rather than logical elements, defaults to `NaN` aware relational comparator + * that evaluates `NaN` as greater than all other values. + * @throw cudf::logic_error if the input table was preprocessed to transform any nested children + * columns into integer columns but `PhysicalElementComparator` is not + * `sorting_physical_element_comparator`. * @param nullate Indicates if any input column contains nulls. * @param comparator Physical element relational comparison functor. * @return A binary callable object. @@ -887,8 +966,10 @@ class self_comparator { template - auto less(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const noexcept + auto less(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const { + d_t->check_physical_element_comparator(); + return less_comparator{ device_row_comparator{ nullate, @@ -906,9 +987,10 @@ class self_comparator { template - auto less_equivalent(Nullate nullate = {}, - PhysicalElementComparator comparator = {}) const noexcept + auto less_equivalent(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const { + d_t->check_physical_element_comparator(); + return less_equivalent_comparator{ device_row_comparator{ nullate, @@ -983,13 +1065,13 @@ class two_table_comparator { * @param left The left table to compare * @param right The right table to compare * @param column_order Optional, host array the same length as a row that indicates the desired - * ascending/descending order of each column in a row. If empty, it is assumed all columns are - * sorted in ascending order. + * ascending/descending order of each column in a row. If empty, it is assumed all columns + * are sorted in ascending order. * @param null_precedence Optional, device array the same length as a row and indicates how null - * values compare to all other for every column. If empty, then null precedence would be - * `null_order::BEFORE` for all columns. + * values compare to all other for every column. If empty, then null precedence would be + * `null_order::BEFORE` for all columns. * @param stream The stream to construct this object on. Not the stream that will be used for - * comparisons using this object. + * comparisons using this object. */ two_table_comparator(table_view const& left, table_view const& right, @@ -1004,6 +1086,10 @@ class two_table_comparator { * This constructor allows independently constructing a `preprocessed_table` and sharing it among * multiple comparators. * + * The preprocessed_table(s) should have been pre-generated together using the factory function + * `preprocessed_table::create(table_view const&, table_view const&)`. Otherwise, the comparison + * results between two tables may be incorrect. + * * @param left A table preprocessed for lexicographic comparison * @param right A table preprocessed for lexicographic comparison */ @@ -1029,9 +1115,9 @@ class two_table_comparator { * `j` of the left table. * * @note The operator overloads in sub-class `element_comparator` are templated via the - * `type_dispatcher` to help select an overload instance for each column in a table. - * So, `cudf::is_nested` will return `true` if the table has nested-type columns, - * but it will be a runtime error if template parameter `has_nested_columns != true`. + * `type_dispatcher` to help select an overload instance for each column in a table. + * So, `cudf::is_nested` will return `true` if the table has nested-type columns, + * but it will be a runtime error if template parameter `has_nested_columns != true`. * * @tparam has_nested_columns compile-time optimization for primitive types. * This template parameter is to be used by the developer by querying @@ -1040,8 +1126,11 @@ class two_table_comparator { * overloads for primitive types. * @tparam Nullate A cudf::nullate type describing whether to check for nulls. * @tparam PhysicalElementComparator A relational comparator functor that compares individual - * values rather than logical elements, defaults to `NaN` aware relational comparator that - * evaluates `NaN` as greater than all other values. + * values rather than logical elements, defaults to `NaN` aware relational comparator + * that evaluates `NaN` as greater than all other values. + * @throw cudf::logic_error if the input tables were preprocessed to transform any nested children + * columns into integer columns but `PhysicalElementComparator` is not + * `sorting_physical_element_comparator`. * @param nullate Indicates if any input column contains nulls. * @param comparator Physical element relational comparison functor. * @return A binary callable object. @@ -1049,8 +1138,11 @@ class two_table_comparator { template - auto less(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const noexcept + auto less(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const { + d_left_table->check_physical_element_comparator(); + d_right_table->check_physical_element_comparator(); + return less_comparator{strong_index_comparator_adapter{ device_row_comparator{ nullate, @@ -1068,9 +1160,11 @@ class two_table_comparator { template - auto less_equivalent(Nullate nullate = {}, - PhysicalElementComparator comparator = {}) const noexcept + auto less_equivalent(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const { + d_left_table->check_physical_element_comparator(); + d_right_table->check_physical_element_comparator(); + return less_equivalent_comparator{strong_index_comparator_adapter{ device_row_comparator{ nullate, diff --git a/cpp/src/table/row_operators.cu b/cpp/src/table/row_operators.cu index 94636f5e6ca..9f3a5bcdfea 100644 --- a/cpp/src/table/row_operators.cu +++ b/cpp/src/table/row_operators.cu @@ -14,8 +14,13 @@ * limitations under the License. */ +#include + #include #include +#include +#include +#include #include #include #include @@ -82,6 +87,12 @@ table_view remove_struct_child_offsets(table_view table) return table_view(cols); } +/** + * @brief The enum to specify whether the `decompose_structs` function will process lists columns + * (at any nested level) or will output them unchanged. + */ +enum class decompose_lists_column : bool { YES, NO }; + /** * @brief Decompose all struct columns in a table * @@ -133,8 +144,13 @@ table_view remove_struct_child_offsets(table_view table) * | | * i f * - * When list columns are present, the decomposition is performed similarly to pure structs but list - * parent columns are NOT pruned + * In the case of structs column with a lists column as its first child such as + * `Struct, float>`, after decomposition we get three columns `Struct<>`, + * `List`, and `float`. + * + * When list columns are present, depending on the input flag `decompose_lists`, the decomposition + * can be performed similarly to pure structs but list parent columns are NOT pruned. The list + * parents are still needed to define the range of elements in the leaf that belong to the same row. * * For example, if the original table has a column `List>`, * @@ -152,20 +168,20 @@ table_view remove_struct_child_offsets(table_view table) * | * i * - * The list parents are still needed to define the range of elements in the leaf that belong to the - * same row. - * - * In the case of structs column with a lists column as its first child such as - * `Struct, float>`, after decomposition we get three columns `Struct<>`, - * `List`, and `float`. + * Note that the `decompose_lists` flag should be specified as follow: + * - `decompose_lists_column::YES` when preprocessing a table for equality comparison. + * - `decompose_lists_column::NO` when preprocessing a table for lexicographic comparison, + * since we need to keep all lists columns intact to input into the next preprocessing step. * * @param table The table whose struct columns to decompose. + * @param decompose_lists Whether to decompose lists columns * @param column_order The per-column order if using output with lexicographic comparison * @param null_precedence The per-column null precedence * @return A tuple containing a table with all struct columns decomposed, new corresponding column * orders and null precedences and depths of the linearized branches */ auto decompose_structs(table_view table, + decompose_lists_column decompose_lists, host_span column_order = {}, host_span null_precedence = {}) { @@ -186,7 +202,7 @@ auto decompose_structs(table_view table, std::vector* branch, int depth) { branch->push_back(c); - if (c->type().id() == type_id::LIST) { + if (decompose_lists == decompose_lists_column::YES && c->type().id() == type_id::LIST) { recursive_child( c->children[lists_column_view::child_column_index].get(), branch, depth + 1); } else if (c->type().id() == type_id::STRUCT) { @@ -317,10 +333,12 @@ void check_lex_compatibility(table_view const& input) if (c.type().id() == type_id::LIST) { auto const& list_col = lists_column_view(c); CUDF_EXPECTS(list_col.child().type().id() != type_id::STRUCT, - "Cannot lexicographic compare a table with a LIST of STRUCT column"); + "Cannot lexicographically compare a table with a LIST of STRUCT column"); check_column(list_col.child()); } else if (c.type().id() == type_id::STRUCT) { for (auto child = c.child_begin(); child < c.child_end(); ++child) { + CUDF_EXPECTS(child->type().id() != type_id::LIST, + "Cannot lexicographically compare a table with a STRUCT of LIST column"); check_column(*child); } } @@ -373,51 +391,391 @@ namespace row { namespace lexicographic { +namespace { + +/** + * @brief Transform any nested lists-of-structs column into lists-of-integers column. + * + * For a lists-of-structs column at any nested level, its child structs column will be replaced by a + * `size_type` column computed as its ranks. + * + * If the input column is not lists-of-structs, or does not contain lists-of-structs at any nested + * level, the input will be passed through without any changes. + * + * @param lhs The input lhs column to transform + * @param rhs The input rhs column to transform (if available) + * @param column_null_order The flag indicating how nulls compare to non-null values + * @param stream CUDA stream used for device memory operations and kernel launches + * @return A tuple consisting of new column_view representing the transformed input, along with + * their ranks column(s) (of `size_type` type) and possibly new list offsets generated + * during the transformation process + */ +std::tuple, + std::vector>, + std::vector>> +transform_lists_of_structs(column_view const& lhs, + std::optional const& rhs_opt, + null_order column_null_order, + rmm::cuda_stream_view stream) +{ + auto const default_mr = rmm::mr::get_current_device_resource(); + + // If the input is not sliced, just replace the input child by new_child. + // Otherwise, we have to generate new offsets and replace both offsets/child of the input by the + // new ones. This is because the new child here is generated by ranking and always has zero + // offset thus cannot replace the input child if it is sliced. + // The new offsets column needs to be returned and kept alive. + auto const replace_child = [&](column_view const& input, + column_view const& new_child, + std::vector>& out_cols) { + auto const make_output = [&input](auto const& offsets_cv, auto const& child_cv) { + return column_view{data_type{type_id::LIST}, + input.size(), + nullptr, + input.null_mask(), + input.null_count(), + 0, + {offsets_cv, child_cv}}; + }; + + if (input.offset() == 0) { + return make_output(input.child(lists_column_view::offsets_column_index), new_child); + } + + out_cols.emplace_back( + cudf::lists::detail::get_normalized_offsets(lists_column_view{input}, stream, default_mr)); + return make_output(out_cols.back()->view(), new_child); + }; + + // Dense ranks should be used instead of first rank. + // Consider this example: `input = [ [{0, "a"}, {3, "c"}], [{0, "a"}, {2, "b"}] ]`. + // If first rank is used, `transformed_input = [ [0, 3], [1, 2] ]`. Comparing them will lead + // to the result row(0) < row(1) which is incorrect. + // With dense rank, `transformed_input = [ [0, 2], [0, 1] ]`, producing correct comparison. + // + // In addition, since the ranked structs column(s) are nested child column instead of + // top-level column, the column order should be fixed to the same values in all situations. + // For example, with the same input above, using the fixed values for column order + // (`order::ASCENDING`), we have `transformed_input = [ [0, 2], [0, 1] ]`. Sorting of + // `transformed_input` will produce the same result as sorting `input` regardless of sorting + // order (ASC or DESC). + auto const compute_ranks = [&](auto const& input) { + return cudf::detail::rank(input, + rank_method::DENSE, + order::ASCENDING, + null_policy::EXCLUDE, + column_null_order, + false /*percentage*/, + stream, + default_mr); + }; + + std::vector> out_cols_lhs; + std::vector> out_cols_rhs; + + if (lhs.type().id() == type_id::LIST) { + auto const child_lhs = cudf::lists_column_view{lhs}.get_sliced_child(stream); + + // Found a lists-of-structs column. + if (child_lhs.type().id() == type_id::STRUCT) { + if (rhs_opt) { // rhs table is available + auto const child_rhs = cudf::lists_column_view{rhs_opt.value()}.get_sliced_child(stream); + auto const concatenated_children = cudf::detail::concatenate( + std::vector{child_lhs, child_rhs}, stream, default_mr); + + auto const ranks = compute_ranks(concatenated_children->view()); + auto const ranks_slices = cudf::detail::slice( + ranks->view(), + {0, child_lhs.size(), child_lhs.size(), child_lhs.size() + child_rhs.size()}, + stream); + + out_cols_lhs.emplace_back(std::make_unique(ranks_slices.front())); + out_cols_rhs.emplace_back(std::make_unique(ranks_slices.back())); + + auto transformed_lhs = replace_child(lhs, out_cols_lhs.back()->view(), out_cols_lhs); + auto transformed_rhs = + replace_child(rhs_opt.value(), out_cols_rhs.back()->view(), out_cols_rhs); + + return {std::move(transformed_lhs), + std::optional{std::move(transformed_rhs)}, + std::move(out_cols_lhs), + std::move(out_cols_rhs)}; + } else { // rhs table is not available + out_cols_lhs.emplace_back(compute_ranks(child_lhs)); + auto transformed_lhs = replace_child(lhs, out_cols_lhs.back()->view(), out_cols_lhs); + + return {std::move(transformed_lhs), + std::nullopt, + std::move(out_cols_lhs), + std::move(out_cols_rhs)}; + } + } + // Found a lists-of-lists column. + else if (child_lhs.type().id() == type_id::LIST) { + auto const child_rhs_opt = + rhs_opt + ? std::optional{cudf::lists_column_view{rhs_opt.value()}.get_sliced_child( + stream)} + : std::nullopt; + + // Recursively call transformation on the child column. + auto [new_child_lhs, new_child_rhs_opt, out_cols_child_lhs, out_cols_child_rhs] = + transform_lists_of_structs(child_lhs, child_rhs_opt, column_null_order, stream); + + // Only transform the current pair of columns if their children have been transformed. + if (out_cols_child_lhs.size() > 0 || out_cols_child_rhs.size() > 0) { + out_cols_lhs.insert(out_cols_lhs.end(), + std::make_move_iterator(out_cols_child_lhs.begin()), + std::make_move_iterator(out_cols_child_lhs.end())); + out_cols_rhs.insert(out_cols_rhs.end(), + std::make_move_iterator(out_cols_child_rhs.begin()), + std::make_move_iterator(out_cols_child_rhs.end())); + + auto transformed_lhs = replace_child(lhs, new_child_lhs, out_cols_lhs); + if (rhs_opt) { + auto transformed_rhs = + replace_child(rhs_opt.value(), new_child_rhs_opt.value(), out_cols_rhs); + + return {std::move(transformed_lhs), + std::optional{std::move(transformed_rhs)}, + std::move(out_cols_lhs), + std::move(out_cols_rhs)}; + } else { + return {std::move(transformed_lhs), + std::nullopt, + std::move(out_cols_lhs), + std::move(out_cols_rhs)}; + } + } + } + // else: child is not STRUCT or LIST: just go to the end of this function, no transformation. + } + // else: lhs.type().id() != type_id::LIST. + // In such situations, lhs.type().id() can still be type_id::STRUCT. However, any + // structs-of-lists should be decomposed into empty struct type `Struct<>` before being + // processed by this function so we do nothing here. + + // Passthrough: nothing changed. + return {lhs, rhs_opt, std::move(out_cols_lhs), std::move(out_cols_rhs)}; +} + +/** + * @brief Transform any nested lists-of-structs column in the given table(s) into lists-of-integers + * column. + * + * If the rhs table is specified, its shape should be pre-checked to match with the shape of lhs + * table using `check_shape_compatibility` before being passed into this function. + * + * @param lhs The input lhs table to transform + * @param rhs The input rhs table to transform (if available) + * @param null_precedence Optional, an array having the same length as the number of columns in + * the input tables that indicates how null values compare to all other. If it is empty, + * the order `null_order::BEFORE` will be used for all columns. + * @param stream CUDA stream used for device memory operations and kernel launches + * @return A tuple consisting of new table_view representing the transformed input, along with + * the ranks columns (of `size_type` type) and possibly new list offsets generated during + * the transformation process + */ +std::tuple, + std::vector>, + std::vector>> +transform_lists_of_structs(table_view const& lhs, + std::optional const& rhs, + host_span null_precedence, + rmm::cuda_stream_view stream) +{ + std::vector transformed_lhs_cvs; + std::vector transformed_rhs_cvs; + std::vector> out_cols_lhs; + std::vector> out_cols_rhs; + + for (size_type col_idx = 0; col_idx < lhs.num_columns(); ++col_idx) { + auto const& lhs_col = lhs.column(col_idx); + auto const rhs_col_opt = + rhs ? std::optional{rhs.value().column(col_idx)} : std::nullopt; + + auto [transformed_lhs, transformed_rhs_opt, curr_out_cols_lhs, curr_out_cols_rhs] = + transform_lists_of_structs( + lhs_col, + rhs_col_opt, + null_precedence.empty() ? null_order::BEFORE : null_precedence[col_idx], + stream); + + transformed_lhs_cvs.emplace_back(std::move(transformed_lhs)); + if (rhs) { transformed_rhs_cvs.emplace_back(std::move(transformed_rhs_opt.value())); } + + out_cols_lhs.insert(out_cols_lhs.end(), + std::make_move_iterator(curr_out_cols_lhs.begin()), + std::make_move_iterator(curr_out_cols_lhs.end())); + out_cols_rhs.insert(out_cols_rhs.end(), + std::make_move_iterator(curr_out_cols_rhs.begin()), + std::make_move_iterator(curr_out_cols_rhs.end())); + } + + return {table_view{transformed_lhs_cvs}, + rhs ? std::optional{table_view{transformed_rhs_cvs}} : std::nullopt, + std::move(out_cols_lhs), + std::move(out_cols_rhs)}; +} + +} // namespace + std::shared_ptr preprocessed_table::create( - table_view const& t, + table_view const& preprocessed_input, + std::vector&& verticalized_col_depths, + std::vector>&& transformed_columns, host_span column_order, host_span null_precedence, + bool has_ranked_children, rmm::cuda_stream_view stream) { - check_lex_compatibility(t); - - auto [verticalized_lhs, new_column_order, new_null_precedence, verticalized_col_depths] = - decompose_structs(t, column_order, null_precedence); + check_lex_compatibility(preprocessed_input); - auto d_t = table_device_view::create(verticalized_lhs, stream); - auto d_column_order = detail::make_device_uvector_async( - new_column_order, stream, rmm::mr::get_current_device_resource()); + auto d_table = table_device_view::create(preprocessed_input, stream); + auto d_column_order = + detail::make_device_uvector_async(column_order, stream, rmm::mr::get_current_device_resource()); auto d_null_precedence = detail::make_device_uvector_async( - new_null_precedence, stream, rmm::mr::get_current_device_resource()); + null_precedence, stream, rmm::mr::get_current_device_resource()); auto d_depths = detail::make_device_uvector_async( verticalized_col_depths, stream, rmm::mr::get_current_device_resource()); - if (detail::has_nested_columns(t)) { - auto [dremel_data, d_dremel_device_view] = list_lex_preprocess(verticalized_lhs, stream); + if (detail::has_nested_columns(preprocessed_input)) { + auto [dremel_data, d_dremel_device_view] = list_lex_preprocess(preprocessed_input, stream); return std::shared_ptr( - new preprocessed_table(std::move(d_t), + new preprocessed_table(std::move(d_table), std::move(d_column_order), std::move(d_null_precedence), std::move(d_depths), std::move(dremel_data), - std::move(d_dremel_device_view))); + std::move(d_dremel_device_view), + std::move(transformed_columns), + has_ranked_children)); } else { - return std::shared_ptr(new preprocessed_table(std::move(d_t), - std::move(d_column_order), - std::move(d_null_precedence), - std::move(d_depths))); + return std::shared_ptr( + new preprocessed_table(std::move(d_table), + std::move(d_column_order), + std::move(d_null_precedence), + std::move(d_depths), + std::move(transformed_columns), + has_ranked_children)); } } +std::shared_ptr preprocessed_table::create( + table_view const& input, + host_span column_order, + host_span null_precedence, + rmm::cuda_stream_view stream) +{ + auto [decomposed_input, new_column_order, new_null_precedence, verticalized_col_depths] = + decompose_structs(input, decompose_lists_column::NO, column_order, null_precedence); + + // Unused variables are generated for rhs table which is not available here. + [[maybe_unused]] auto [transformed_input, unused_0, transformed_columns, unused_1] = + transform_lists_of_structs(decomposed_input, std::nullopt, new_null_precedence, stream); + + auto const has_ranked_children = !transformed_columns.empty(); + return create(transformed_input, + std::move(verticalized_col_depths), + std::move(transformed_columns), + new_column_order, + new_null_precedence, + has_ranked_children, + stream); +} + +std::pair, std::shared_ptr> +preprocessed_table::create(table_view const& lhs, + table_view const& rhs, + host_span column_order, + host_span null_precedence, + rmm::cuda_stream_view stream) +{ + check_shape_compatibility(lhs, rhs); + + auto [decomposed_lhs, + new_column_order_lhs, + new_null_precedence_lhs, + verticalized_col_depths_lhs] = + decompose_structs(lhs, decompose_lists_column::NO, column_order, null_precedence); + + // Unused variables are new column order and null order for rhs, which are the same as for lhs + // so we don't need them. + [[maybe_unused]] auto [decomposed_rhs, unused0, unused1, verticalized_col_depths_rhs] = + decompose_structs(rhs, decompose_lists_column::NO, column_order, null_precedence); + + // Transform any (nested) lists-of-structs column into lists-of-integers column. + auto [transformed_lhs, transformed_rhs_opt, transformed_columns_lhs, transformed_columns_rhs] = + transform_lists_of_structs(decomposed_lhs, decomposed_rhs, new_null_precedence_lhs, stream); + + // This should be the same for both lhs and rhs but not all the time, such as when one table + // has 0 rows while the other has >0 rows. So we check separately for each of them. + auto const has_ranked_children_lhs = !transformed_columns_lhs.empty(); + auto const has_ranked_children_rhs = !transformed_columns_rhs.empty(); + + return {create(transformed_lhs, + std::move(verticalized_col_depths_lhs), + std::move(transformed_columns_lhs), + new_column_order_lhs, + new_null_precedence_lhs, + has_ranked_children_lhs, + stream), + create(transformed_rhs_opt.value(), + std::move(verticalized_col_depths_rhs), + std::move(transformed_columns_rhs), + new_column_order_lhs, + new_null_precedence_lhs, + has_ranked_children_rhs, + stream)}; +} + +preprocessed_table::preprocessed_table( + table_device_view_owner&& table, + rmm::device_uvector&& column_order, + rmm::device_uvector&& null_precedence, + rmm::device_uvector&& depths, + std::vector&& dremel_data, + rmm::device_uvector&& dremel_device_views, + std::vector>&& transformed_columns, + bool has_ranked_children) + : _t(std::move(table)), + _column_order(std::move(column_order)), + _null_precedence(std::move(null_precedence)), + _depths(std::move(depths)), + _dremel_data(std::move(dremel_data)), + _dremel_device_views(std::move(dremel_device_views)), + _transformed_columns(std::move(transformed_columns)), + _has_ranked_children(has_ranked_children) +{ +} + +preprocessed_table::preprocessed_table(table_device_view_owner&& table, + rmm::device_uvector&& column_order, + rmm::device_uvector&& null_precedence, + rmm::device_uvector&& depths, + std::vector>&& transformed_columns, + bool has_ranked_children) + : _t(std::move(table)), + _column_order(std::move(column_order)), + _null_precedence(std::move(null_precedence)), + _depths(std::move(depths)), + _dremel_data{}, + _dremel_device_views{}, + _transformed_columns(std::move(transformed_columns)), + _has_ranked_children(has_ranked_children) +{ +} + two_table_comparator::two_table_comparator(table_view const& left, table_view const& right, host_span column_order, host_span null_precedence, rmm::cuda_stream_view stream) - : d_left_table{preprocessed_table::create(left, column_order, null_precedence, stream)}, - d_right_table{preprocessed_table::create(right, column_order, null_precedence, stream)} { - check_shape_compatibility(left, right); + std::tie(d_left_table, d_right_table) = + preprocessed_table::create(left, right, column_order, null_precedence, stream); } } // namespace lexicographic @@ -432,7 +790,8 @@ std::shared_ptr preprocessed_table::create(table_view const& auto [null_pushed_table, nullable_data] = structs::detail::push_down_nulls(t, stream, rmm::mr::get_current_device_resource()); auto struct_offset_removed_table = remove_struct_child_offsets(null_pushed_table); - auto verticalized_t = std::get<0>(decompose_structs(struct_offset_removed_table)); + auto verticalized_t = + std::get<0>(decompose_structs(struct_offset_removed_table, decompose_lists_column::YES)); auto d_t = table_device_view_owner(table_device_view::create(verticalized_t, stream)); return std::shared_ptr(new preprocessed_table( diff --git a/cpp/tests/search/search_list_test.cpp b/cpp/tests/search/search_list_test.cpp index ca17858a273..48711c21715 100644 --- a/cpp/tests/search/search_list_test.cpp +++ b/cpp/tests/search/search_list_test.cpp @@ -34,9 +34,8 @@ using structs_col = cudf::test::structs_column_wrapper; using strings_col = cudf::test::strings_column_wrapper; constexpr cudf::test::debug_output_level verbosity{cudf::test::debug_output_level::FIRST_ERROR}; -constexpr int32_t null{0}; // Mark for null child elements at the current level -constexpr int32_t XXX{0}; // Mark for null elements at all levels -constexpr int32_t dont_care{0}; // Mark for elements that will be sliced off +constexpr int32_t null{0}; // Mark for null child elements at the current level +constexpr int32_t XXX{0}; // Mark for null elements at all levels using TestTypes = cudf::test::Concat; using lists_col = cudf::test::lists_column_wrapper; - auto const haystack_original = - lists_col{{dont_care, dont_care}, {dont_care}, {1, 2}, {1}, {}, {1, 3}, {dont_care, dont_care}}; - auto const haystack = cudf::slice(haystack_original, {2, 6})[0]; + auto const haystack_original = lists_col{{0, 0}, {0}, {1, 2}, {1}, {}, {1, 3}, {0, 0}}; + auto const haystack = cudf::slice(haystack_original, {2, 6})[0]; auto const needle1 = [] { auto child = tdata_col{1, 2}; @@ -107,7 +105,7 @@ TYPED_TEST(TypedListsContainsTestScalarNeedle, SlicedColumnInput) return cudf::list_scalar(child); }(); auto const needle3 = [] { - auto child = tdata_col{dont_care, dont_care}; + auto child = tdata_col{0, 0}; return cudf::list_scalar(child); }(); @@ -187,8 +185,8 @@ TYPED_TEST(TypedListsContainsTestScalarNeedle, SlicedInputHavingNulls) using tdata_col = cudf::test::fixed_width_column_wrapper; using lists_col = cudf::test::lists_column_wrapper; - auto const haystack_original = lists_col{{{dont_care, dont_care}, - {dont_care} /*NULL*/, + auto const haystack_original = lists_col{{{0, 0}, + {0} /*NULL*/, lists_col{{1, null}, null_at(1)}, {1}, {} /*NULL*/, @@ -196,7 +194,7 @@ TYPED_TEST(TypedListsContainsTestScalarNeedle, SlicedInputHavingNulls) {4}, {} /*NULL*/, {1, 1}, - {dont_care}}, + {0}}, nulls_at({1, 4, 7})}; auto const haystack = cudf::slice(haystack_original, {2, 9})[0]; @@ -209,7 +207,7 @@ TYPED_TEST(TypedListsContainsTestScalarNeedle, SlicedInputHavingNulls) return cudf::list_scalar(child); }(); auto const needle3 = [] { - auto child = tdata_col{dont_care, dont_care}; + auto child = tdata_col{0, 0}; return cudf::list_scalar(child); }(); @@ -250,13 +248,12 @@ TYPED_TEST(TypedListContainsTestColumnNeedles, SlicedInputNoNulls) { using lists_col = cudf::test::lists_column_wrapper; - auto const haystack_original = lists_col{ - {dont_care, dont_care}, {dont_care}, {0, 1}, {2}, {3, 4, 5}, {2, 3, 4}, {}, {0, 2, 0}}; + auto const haystack_original = + lists_col{{0, 0}, {0}, {0, 1}, {2}, {3, 4, 5}, {2, 3, 4}, {}, {0, 2, 0}}; auto const haystack = cudf::slice(haystack_original, {2, 8})[0]; - auto const needles_original = - lists_col{{dont_care}, {0, 1}, {0, 0}, {3, 5, 4}, {}, {dont_care, dont_care}, {} /*dont_care*/}; - auto const needles = cudf::slice(needles_original, {1, 5})[0]; + auto const needles_original = lists_col{{0}, {0, 1}, {0, 0}, {3, 5, 4}, {}, {0, 0}, {} /*0*/}; + auto const needles = cudf::slice(needles_original, {1, 5})[0]; auto const expected = bools_col{1, 0, 0, 1}; auto const result = cudf::contains(haystack, needles); @@ -267,8 +264,8 @@ TYPED_TEST(TypedListContainsTestColumnNeedles, SlicedInputHavingNulls) { using lists_col = cudf::test::lists_column_wrapper; - auto const haystack_original = lists_col{{{dont_care, dont_care}, - {dont_care} /*NULL*/, + auto const haystack_original = lists_col{{{0, 0}, + {0} /*NULL*/, lists_col{{1, null}, null_at(1)}, {1}, {} /*NULL*/, @@ -276,12 +273,12 @@ TYPED_TEST(TypedListContainsTestColumnNeedles, SlicedInputHavingNulls) {4}, {} /*NULL*/, {1, 1}, - {dont_care}}, + {0}}, nulls_at({1, 4, 7})}; auto const haystack = cudf::slice(haystack_original, {2, 9})[0]; - auto const needles_original = lists_col{{{dont_care, dont_care}, - {dont_care} /*NULL*/, + auto const needles_original = lists_col{{{0, 0}, + {0} /*NULL*/, lists_col{{1, null}, null_at(1)}, {1}, {} /*NULL*/, @@ -289,7 +286,7 @@ TYPED_TEST(TypedListContainsTestColumnNeedles, SlicedInputHavingNulls) {4}, {} /*NULL*/, {}, - {dont_care}}, + {0}}, nulls_at({1, 4, 7})}; auto const needles = cudf::slice(needles_original, {2, 9})[0]; @@ -347,9 +344,20 @@ TYPED_TEST(TypedListContainsTestColumnNeedles, ListsOfStructs) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, *result, verbosity); } -struct ListLowerBound : public cudf::test::BaseFixture {}; +auto search_bounds(cudf::table_view const& t, + cudf::table_view const& values, + std::vector const& column_order, + std::vector const& null_precedence = { + cudf::null_order::BEFORE}) +{ + auto result_lower_bound = cudf::lower_bound(t, values, column_order, null_precedence); + auto result_upper_bound = cudf::upper_bound(t, values, column_order, null_precedence); + return std::pair(std::move(result_lower_bound), std::move(result_upper_bound)); +} + +struct ListBinarySearch : public cudf::test::BaseFixture {}; -TEST_F(ListLowerBound, ListWithNulls) +TEST_F(ListBinarySearch, ListWithNulls) { { using lcw = cudf::test::lists_column_wrapper; @@ -360,29 +368,31 @@ TEST_F(ListLowerBound, ListWithNulls) }; auto const needles = lcw{ - lcw{{0, 4.22671e+32}, null_at(0)}, + lcw{{null, 4.22671e+32}, null_at(0)}, }; - auto const expect = int32s_col{0}; - auto const result = cudf::lower_bound(cudf::table_view{{haystack}}, - cudf::table_view{{needles}}, - {cudf::order::ASCENDING}, - {cudf::null_order::BEFORE}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, *result); + auto const expected = int32s_col{0}; + auto const [result_lower_bound, result_upper_bound] = + search_bounds(cudf::table_view{{haystack}}, + cudf::table_view{{needles}}, + {cudf::order::ASCENDING}, + {cudf::null_order::BEFORE}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, *result_lower_bound, verbosity); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, *result_upper_bound, verbosity); } { using lcw = cudf::test::lists_column_wrapper; auto const col1 = lcw{ - lcw{{0}, null_at(0)}, // 0 - lcw{-80}, // 1 - lcw{-17}, // 2 + lcw{{null}, null_at(0)}, // 0 + lcw{-80}, // 1 + lcw{-17}, // 2 }; auto const col2 = lcw{ - lcw{27}, // 0 - lcw{{0}, null_at(0)}, // 1 - lcw{}, // 2 + lcw{27}, // 0 + lcw{{null}, null_at(0)}, // 1 + lcw{}, // 2 }; auto const val1 = lcw{ @@ -399,8 +409,262 @@ TEST_F(ListLowerBound, ListWithNulls) std::vector null_order_flags{cudf::null_order::BEFORE, cudf::null_order::BEFORE}; - auto const expect = int32s_col{3}; - auto const result = cudf::lower_bound(input, values, column_order, null_order_flags); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, *result); + auto const expected = int32s_col{3}; + auto const [result_lower_bound, result_upper_bound] = search_bounds( + cudf::table_view{{input}}, cudf::table_view{{values}}, column_order, null_order_flags); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, *result_lower_bound, verbosity); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, *result_upper_bound, verbosity); } } + +TEST_F(ListBinarySearch, ListsOfStructs) +{ + // Haystack must be pre-sorted. + auto const haystack = [] { + auto offsets = int32s_col{0, 2, 3, 4, 5, 7, 10, 13, 16, 18}; + // clang-format off + auto data1 = int32s_col{1, 2, + 3, + 3, + 3, + 4, 5, + 4, 5, 4, + 4, 5, 4, + 4, 5, 4, + 4, 6 + }; + auto data2 = int32s_col{1, 2, + 3, + 3, + 3, + 4, 5, + 4, 5, 4, + 4, 5, 4, + 4, 5, 4, + 5, 1 + }; + // clang-format on + auto child = structs_col{{data1, data2}}; + return cudf::make_lists_column(9, offsets.release(), child.release(), 0, {}); + }(); + + auto const needles = [] { + auto offsets = int32s_col{0, 3, 4, 6, 8, 10, 13, 14, 15}; + // clang-format off + auto data1 = int32s_col{1, 2, 1, + 3, + 4, 1, + 0, 1, + 1, 0, + 1, 3, 5, + 3, + 3 + }; + auto data2 = int32s_col{1, 3, 0, + 3, + 1, 2, + 1, 1, + 1, 2, + 0, 2, 2, + 3, + 3 + }; + // clang-format on + auto child = structs_col{{data1, data2}}; + return cudf::make_lists_column(8, offsets.release(), child.release(), 0, {}); + }(); + + auto const [result_lower_bound, result_upper_bound] = search_bounds( + cudf::table_view{{*haystack}}, cudf::table_view{{*needles}}, {cudf::order::ASCENDING}); + auto const expected_lower_bound = int32s_col{1, 1, 4, 0, 0, 0, 1, 1}; + auto const expected_upper_bound = int32s_col{1, 4, 4, 0, 0, 0, 4, 4}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_lower_bound, *result_lower_bound, verbosity); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_upper_bound, *result_upper_bound, verbosity); +} + +TEST_F(ListBinarySearch, ListsOfEqualStructsInTwoTables) +{ + // Haystack must be pre-sorted. + auto const haystack = [] { + auto offsets = int32s_col{0, 2, 3, 4, 5, 7, 10, 13, 16, 18}; + // clang-format off + auto data1 = int32s_col{1, 2, + 3, + 3, + 3, + 4, 5, + 4, 5, 4, + 4, 5, 4, + 4, 5, 4, + 4, 6 + }; + auto data2 = int32s_col{1, 2, + 3, + 3, + 3, + 4, 5, + 4, 5, 4, + 4, 5, 4, + 4, 5, 4, + 5, 1 + }; + // clang-format on + auto child = structs_col{{data1, data2}}; + return cudf::make_lists_column(9, offsets.release(), child.release(), 0, {}); + }(); + + auto const needles = [] { + auto offsets = int32s_col{0, 2, 3, 4, 5, 7, 10, 13, 15, 17}; + // clang-format off + auto data1 = int32s_col{1, 2, + 3, + 4, + 5, + 4, 5, + 5, 5, 4, + 4, 5, 4, + 4, 4, + 4, 6 + }; + auto data2 = int32s_col{1, 2, + 3, + 4, + 5, + 4, 5, + 5, 5, 4, + 4, 5, 4, + 4, 4, + 5, 1 + }; + // clang-format on + auto child = structs_col{{data1, data2}}; + return cudf::make_lists_column(9, offsets.release(), child.release(), 0, {}); + }(); + + // In this search, the two table have many equal structs. + // This help to verify the internal implementation of two-table lex comparator in which the + // structs column of two input tables are concatenated, ranked, then split. + auto const [result_lower_bound, result_upper_bound] = search_bounds( + cudf::table_view{{*haystack}}, cudf::table_view{{*needles}}, {cudf::order::ASCENDING}); + auto const expected_lower_bound = int32s_col{0, 1, 4, 9, 4, 9, 5, 4, 8}; + auto const expected_upper_bound = int32s_col{1, 4, 4, 9, 5, 9, 8, 4, 9}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_lower_bound, *result_lower_bound, verbosity); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_upper_bound, *result_upper_bound, verbosity); +} + +TEST_F(ListBinarySearch, CrazyListTest) +{ + // Data type: List>>>>> + + // Haystack must be pre-sorted. + auto const haystack = [] { + auto lists_of_structs_of_ints = [] { + auto offsets = int32s_col{0, 2, 3, 4, 5, 7, 10, 13, 16, 18}; + // clang-format off + auto data1 = int32s_col{1, 2, + 3, + 3, + 3, + 4, 5, + 4, 5, 4, + 4, 5, 4, + 4, 5, 4, + 4, 6 + }; + auto data2 = int32s_col{1, 2, + 3, + 3, + 3, + 4, 5, + 4, 5, 4, + 4, 5, 4, + 4, 5, 4, + 5, 1 + }; + // clang-format on + auto child = structs_col{{data1, data2}}; + return cudf::make_lists_column(9, offsets.release(), child.release(), 0, {}); + }(); + + auto struct_nested0 = [&] { + std::vector> child_columns; + child_columns.emplace_back(std::move(lists_of_structs_of_ints)); + return cudf::make_structs_column(9, std::move(child_columns), 0, {}); + }(); + + auto struct_nested1 = [&] { + std::vector> child_columns; + child_columns.emplace_back(std::move(struct_nested0)); + return cudf::make_structs_column(9, std::move(child_columns), 0, {}); + }(); + + auto list_nested0 = [&] { + auto offsets = int32s_col{0, 3, 3, 4, 6, 9}; + return cudf::make_lists_column(5, offsets.release(), std::move(struct_nested1), 0, {}); + }(); + + auto offsets = int32s_col{0, 0, 2, 4, 5, 5}; + return cudf::make_lists_column(5, offsets.release(), std::move(list_nested0), 0, {}); + }(); + + auto const needles = [] { + auto lists_of_structs_of_ints = [] { + auto offsets = int32s_col{0, 2, 3, 4, 5, 7, 10, 13, 15, 17}; + // clang-format off + auto data1 = int32s_col{1, 2, + 3, + 4, + 5, + 4, 5, + 5, 5, 4, + 4, 5, 4, + 4, 4, + 4, 6 + }; + auto data2 = int32s_col{1, 2, + 3, + 4, + 5, + 4, 5, + 5, 5, 4, + 4, 5, 4, + 4, 4, + 5, 1 + }; + // clang-format on + auto child = structs_col{{data1, data2}}; + return cudf::make_lists_column(9, offsets.release(), child.release(), 0, {}); + }(); + + auto struct_nested0 = [&] { + std::vector> child_columns; + child_columns.emplace_back(std::move(lists_of_structs_of_ints)); + return cudf::make_structs_column(9, std::move(child_columns), 0, {}); + }(); + + auto struct_nested1 = [&] { + std::vector> child_columns; + child_columns.emplace_back(std::move(struct_nested0)); + return cudf::make_structs_column(9, std::move(child_columns), 0, {}); + }(); + + auto list_nested0 = [&] { + auto offsets = int32s_col{0, 3, 3, 4, 6, 9}; + return cudf::make_lists_column(5, offsets.release(), std::move(struct_nested1), 0, {}); + }(); + + auto offsets = int32s_col{0, 2, 2, 4, 4, 5}; + return cudf::make_lists_column(5, offsets.release(), std::move(list_nested0), 0, {}); + }(); + + // In this search, the two table have many equal structs. + // This help to verify the internal implementation of two-table lex comparator in which the + // structs column of two input tables are concatenated, ranked, then split. + auto const [result_lower_bound, result_upper_bound] = search_bounds( + cudf::table_view{{*haystack}}, cudf::table_view{{*needles}}, {cudf::order::ASCENDING}); + + auto const expected_lower_bound = int32s_col{2, 0, 5, 0, 5}; + auto const expected_upper_bound = int32s_col{2, 1, 5, 1, 5}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_lower_bound, *result_lower_bound, verbosity); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_upper_bound, *result_upper_bound, verbosity); +} diff --git a/cpp/tests/sort/sort_nested_types_tests.cpp b/cpp/tests/sort/sort_nested_types_tests.cpp index 985793a3d91..ff6256c2408 100644 --- a/cpp/tests/sort/sort_nested_types_tests.cpp +++ b/cpp/tests/sort/sort_nested_types_tests.cpp @@ -242,3 +242,202 @@ TEST_F(NestedStructTest, StructsOfStructsHaveListsWithNulls) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); } } + +TEST_F(NestedStructTest, SimpleStructsOfListsOfStructsNoNulls) +{ + auto const input = [] { + auto const make_lists_of_structs = [] { + auto const get_structs = [] { + auto child0 = int32s_col{3, 2, 3, 3, 4, 2, 4, 4, 1, 0, 3, 0, 2, 5, 4}; + auto child1 = int32s_col{0, 4, 3, 2, 1, 1, 5, 1, 5, 5, 4, 2, 4, 1, 3}; + return structs_col{{child0, child1}}; + }; + return cudf::make_lists_column( + 8, int32s_col{0, 3, 5, 6, 6, 8, 10, 12, 15}.release(), get_structs().release(), 0, {}); + }; + + std::vector> children; + children.emplace_back(make_lists_of_structs()); + children.emplace_back(make_lists_of_structs()); + + return cudf::make_structs_column(8, std::move(children), 0, {}); + }(); + + { + auto const expected_order = int32s_col{3, 5, 2, 7, 0, 1, 6, 4}; + auto const order = cudf::stable_sorted_order(cudf::table_view{{*input}}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); + } + + { + auto const expected_order = int32s_col{4, 6, 1, 0, 7, 2, 5, 3}; + auto const order = + cudf::stable_sorted_order(cudf::table_view{{*input}}, {cudf::order::DESCENDING}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); + } +} + +struct NestedListTest : public cudf::test::BaseFixture {}; + +TEST_F(NestedListTest, SimpleListsOfStructsNoNulls) +{ + auto const input = [] { + auto const get_structs = [] { + auto child0 = int32s_col{3, 2, 3, 3, 4, 2, 4, 4, 1, 0, 3, 0, 2, 5, 4}; + auto child1 = int32s_col{0, 4, 3, 2, 1, 1, 5, 1, 5, 5, 4, 2, 4, 1, 3}; + return structs_col{{child0, child1}}; + }; + return cudf::make_lists_column( + 8, int32s_col{0, 3, 5, 6, 6, 8, 10, 12, 15}.release(), get_structs().release(), 0, {}); + }(); + + { + auto const expected_order = int32s_col{3, 5, 2, 7, 0, 1, 6, 4}; + auto const order = cudf::stable_sorted_order(cudf::table_view{{*input}}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); + } + + { + auto const expected_order = int32s_col{4, 6, 1, 0, 7, 2, 5, 3}; + auto const order = + cudf::stable_sorted_order(cudf::table_view{{*input}}, {cudf::order::DESCENDING}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); + } +} + +TEST_F(NestedListTest, SlicedListsOfStructsNoNulls) +{ + auto const input_original = [] { + auto const get_structs = [] { + auto child0 = int32s_col{0, 0, 3, 2, 3, 3, 4, 2, 4, 4, 1, 0, 3, 0, 2, 5, 4, 0}; + auto child1 = int32s_col{0, 0, 0, 4, 3, 2, 1, 1, 5, 1, 5, 5, 4, 2, 4, 1, 3, 0}; + return structs_col{{child0, child1}}; + }; + return cudf::make_lists_column(11, + int32s_col{0, 1, 2, 5, 7, 8, 8, 10, 12, 14, 17, 18}.release(), + get_structs().release(), + 0, + {}); + }(); + auto const input = cudf::slice(*input_original, {2, 10})[0]; + + { + auto const expected_order = int32s_col{3, 5, 2, 7, 0, 1, 6, 4}; + auto const order = cudf::stable_sorted_order(cudf::table_view{{input}}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); + } + + { + auto const expected_order = int32s_col{4, 6, 1, 0, 7, 2, 5, 3}; + auto const order = + cudf::stable_sorted_order(cudf::table_view{{input}}, {cudf::order::DESCENDING}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); + } +} + +TEST_F(NestedListTest, ListsOfEqualStructsNoNulls) +{ + auto const input = [] { + auto const get_structs = [] { + auto child0 = int32s_col{0, 3, 0, 1}; + auto child1 = strings_col{"a", "c", "a", "b"}; + return structs_col{{child0, child1}}; + }; + return cudf::make_lists_column( + 2, int32s_col{0, 2, 4}.release(), get_structs().release(), 0, {}); + }(); + + { + auto const expected_order = int32s_col{1, 0}; + auto const order = cudf::sorted_order(cudf::table_view{{*input}}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); + } + + { + auto const expected_order = int32s_col{0, 1}; + auto const order = cudf::sorted_order(cudf::table_view{{*input}}, {cudf::order::DESCENDING}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); + } +} + +TEST_F(NestedListTest, SimpleListsOfStructsWithNulls) +{ + // [ {null, 2}, {null, null}, {1, 2} ] | 0 + // [] | 1 + // [ {null, null}, {4, 2} ] | 2 + // [] | 3 + // [ {3, 5}, {null, 4} ] | 4 + // [] | 5 + // [ {5, 3}, {5, 0}, {1, 1} ] | 6 + // [ {null, 3}, {5, 2}, {4, 2} ] | 7 + auto const input = [] { + auto const get_structs = [] { + auto child0 = int32s_col{{null, null, 1, null, 4, 3, null, 5, 5, 1, null, 5, 4}, + nulls_at({0, 1, 3, 6, 10})}; + auto child1 = int32s_col{{2, null, 2, null, 2, 5, 4, 3, 0, 1, 3, 2, 2}, nulls_at({1, 3})}; + return structs_col{{child0, child1}, nulls_at({1, 3})}; + }; + return cudf::make_lists_column( + 8, int32s_col{0, 3, 3, 5, 5, 7, 7, 10, 13}.release(), get_structs().release(), 0, {}); + }(); + + { + auto const expected_order = int32s_col{1, 3, 5, 2, 0, 7, 4, 6}; + auto const order = cudf::stable_sorted_order( + cudf::table_view{{*input}}, {cudf::order::ASCENDING}, {cudf::null_order::BEFORE}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); + } + + { + auto const expected_order = int32s_col{6, 4, 7, 0, 2, 1, 3, 5}; + auto const order = cudf::stable_sorted_order( + cudf::table_view{{*input}}, {cudf::order::DESCENDING}, {cudf::null_order::BEFORE}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); + } + + { + auto const expected_order = int32s_col{1, 3, 5, 2, 4, 6, 0, 7}; + auto const order = cudf::stable_sorted_order( + cudf::table_view{{*input}}, {cudf::order::ASCENDING}, {cudf::null_order::AFTER}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); + } + + { + auto const expected_order = int32s_col{7, 0, 6, 4, 2, 1, 3, 5}; + auto const order = cudf::stable_sorted_order( + cudf::table_view{{*input}}, {cudf::order::DESCENDING}, {cudf::null_order::AFTER}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); + } +} + +TEST_F(NestedListTest, ListsOfListsOfStructsNoNulls) +{ + auto const input = [] { + auto const get_structs = [] { + auto child0 = int32s_col{0, 7, 4, 9, 2, 9, 4, 1, 5, 5, 3, 7, 0, 6, 3, 1, 9}; + auto child1 = int32s_col{4, 6, 7, 3, 1, 2, 1, 10, 7, 9, 8, 7, 1, 10, 5, 3, 3}; + return structs_col{{child0, child1}}; + }; + auto lists_of_structs = + cudf::make_lists_column(13, + int32s_col{0, 1, 3, 4, 5, 7, 9, 10, 12, 12, 14, 15, 17, 17}.release(), + get_structs().release(), + 0, + {}); + return cudf::make_lists_column( + 8, int32s_col{0, 3, 4, 6, 6, 8, 10, 10, 13}.release(), std::move(lists_of_structs), 0, {}); + }(); + + { + auto const expected_order = int32s_col{3, 6, 5, 0, 1, 7, 4, 2}; + auto const order = cudf::stable_sorted_order(cudf::table_view{{*input}}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); + } + + { + auto const expected_order = int32s_col{2, 4, 7, 1, 0, 5, 3, 6}; + auto const order = + cudf::stable_sorted_order(cudf::table_view{{*input}}, {cudf::order::DESCENDING}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected_order, order->view()); + } +} diff --git a/cpp/tests/table/experimental_row_operator_tests.cu b/cpp/tests/table/experimental_row_operator_tests.cu index 42f1753494f..5ae1c7d9729 100644 --- a/cpp/tests/table/experimental_row_operator_tests.cu +++ b/cpp/tests/table/experimental_row_operator_tests.cu @@ -19,6 +19,7 @@ #include #include +#include #include #include #include @@ -29,6 +30,8 @@ #include #include +#include +#include #include #include @@ -184,6 +187,34 @@ auto two_table_equality(cudf::table_view lhs, return output; } +template +auto sorted_order( + std::shared_ptr preprocessed_input, + cudf::size_type num_rows, + bool has_nested, + PhysicalElementComparator comparator, + rmm::cuda_stream_view stream) +{ + auto output = cudf::make_numeric_column(cudf::data_type(cudf::type_to_id()), + num_rows, + cudf::mask_state::UNALLOCATED, + stream); + auto const out_begin = output->mutable_view().begin(); + thrust::sequence(rmm::exec_policy(stream), out_begin, out_begin + num_rows, 0); + + auto const table_comparator = + cudf::experimental::row::lexicographic::self_comparator{preprocessed_input}; + if (has_nested) { + auto const comp = table_comparator.less(cudf::nullate::NO{}, comparator); + thrust::stable_sort(rmm::exec_policy(stream), out_begin, out_begin + num_rows, comp); + } else { + auto const comp = table_comparator.less(cudf::nullate::NO{}, comparator); + thrust::stable_sort(rmm::exec_policy(stream), out_begin, out_begin + num_rows, comp); + } + + return output; +} + TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorTwoTables) { using T = TypeParam; @@ -229,6 +260,148 @@ TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorSameTable) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, sorting_got->view()); } +TYPED_TEST(TypedTableViewTest, TestSortSameTableFromTwoTables) +{ + using data_col = cudf::test::fixed_width_column_wrapper; + using int32s_col = cudf::test::fixed_width_column_wrapper; + + auto const col1 = data_col{5, 2, 7, 1, 3}; + auto const col2 = data_col{}; // empty + auto const lhs = cudf::table_view{{col1}}; + auto const empty_rhs = cudf::table_view{{col2}}; + + auto const stream = cudf::get_default_stream(); + auto const test_sort = [stream](auto const& preprocessed, + auto const& input, + auto const& comparator, + auto const& expected) { + auto const order = sorted_order( + preprocessed, input.num_rows(), cudf::detail::has_nested_columns(input), comparator, stream); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, order->view()); + }; + + auto const test_sort_two_tables = [&](auto const& preprocessed_lhs, + auto const& preprocessed_empty_rhs) { + auto const expected_lhs = int32s_col{3, 1, 4, 0, 2}; + test_sort(preprocessed_lhs, + lhs, + cudf::experimental::row::lexicographic::physical_element_comparator{}, + expected_lhs); + test_sort(preprocessed_lhs, + lhs, + cudf::experimental::row::lexicographic::sorting_physical_element_comparator{}, + expected_lhs); + + auto const expected_empty_rhs = int32s_col{}; + test_sort(preprocessed_empty_rhs, + empty_rhs, + cudf::experimental::row::lexicographic::physical_element_comparator{}, + expected_empty_rhs); + test_sort(preprocessed_empty_rhs, + empty_rhs, + cudf::experimental::row::lexicographic::sorting_physical_element_comparator{}, + expected_empty_rhs); + }; + + // Generate preprocessed data for both lhs and lhs at the same time. + // Switching order of lhs and rhs tables then sorting them using their preprocessed data should + // produce exactly the same result. + { + auto const [preprocessed_lhs, preprocessed_empty_rhs] = + cudf::experimental::row::lexicographic::preprocessed_table::create( + lhs, empty_rhs, std::vector{cudf::order::ASCENDING}, {}, stream); + test_sort_two_tables(preprocessed_lhs, preprocessed_empty_rhs); + } + { + auto const [preprocessed_empty_rhs, preprocessed_lhs] = + cudf::experimental::row::lexicographic::preprocessed_table::create( + empty_rhs, lhs, std::vector{cudf::order::ASCENDING}, {}, stream); + test_sort_two_tables(preprocessed_lhs, preprocessed_empty_rhs); + } +} + +TYPED_TEST(TypedTableViewTest, TestSortSameTableFromTwoTablesWithListsOfStructs) +{ + using data_col = cudf::test::fixed_width_column_wrapper; + using int32s_col = cudf::test::fixed_width_column_wrapper; + using strings_col = cudf::test::strings_column_wrapper; + using structs_col = cudf::test::structs_column_wrapper; + + auto const col1 = [] { + auto const get_structs = [] { + auto child0 = data_col{0, 3, 0, 2}; + auto child1 = strings_col{"a", "c", "a", "b"}; + return structs_col{{child0, child1}}; + }; + return cudf::make_lists_column( + 2, int32s_col{0, 2, 4}.release(), get_structs().release(), 0, {}); + }(); + auto const col2 = [] { + auto const get_structs = [] { + auto child0 = data_col{}; + auto child1 = strings_col{}; + return structs_col{{child0, child1}}; + }; + return cudf::make_lists_column(0, int32s_col{}.release(), get_structs().release(), 0, {}); + }(); + + auto const column_order = std::vector{cudf::order::ASCENDING}; + auto const lhs = cudf::table_view{{*col1}}; + auto const empty_rhs = cudf::table_view{{*col2}}; + + auto const stream = cudf::get_default_stream(); + auto const test_sort = [stream](auto const& preprocessed, + auto const& input, + auto const& comparator, + auto const& expected) { + auto const order = sorted_order( + preprocessed, input.num_rows(), cudf::detail::has_nested_columns(input), comparator, stream); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, order->view()); + }; + + auto const test_sort_two_tables = [&](auto const& preprocessed_lhs, + auto const& preprocessed_empty_rhs) { + auto const expected_lhs = int32s_col{1, 0}; + test_sort(preprocessed_lhs, + lhs, + cudf::experimental::row::lexicographic::sorting_physical_element_comparator{}, + expected_lhs); + + auto const expected_empty_rhs = int32s_col{}; + test_sort(preprocessed_empty_rhs, + empty_rhs, + cudf::experimental::row::lexicographic::sorting_physical_element_comparator{}, + expected_empty_rhs); + + EXPECT_THROW(test_sort(preprocessed_lhs, + lhs, + cudf::experimental::row::lexicographic::physical_element_comparator{}, + expected_lhs), + cudf::logic_error); + EXPECT_THROW(test_sort(preprocessed_empty_rhs, + empty_rhs, + cudf::experimental::row::lexicographic::physical_element_comparator{}, + expected_empty_rhs), + cudf::logic_error); + }; + + // Generate preprocessed data for both lhs and lhs at the same time. + // Switching order of lhs and rhs tables then sorting them using their preprocessed data should + // produce exactly the same result. + { + auto const [preprocessed_lhs, preprocessed_empty_rhs] = + cudf::experimental::row::lexicographic::preprocessed_table::create( + lhs, empty_rhs, std::vector{cudf::order::ASCENDING}, {}, stream); + test_sort_two_tables(preprocessed_lhs, preprocessed_empty_rhs); + } + { + auto const [preprocessed_empty_rhs, preprocessed_lhs] = + cudf::experimental::row::lexicographic::preprocessed_table::create( + empty_rhs, lhs, std::vector{cudf::order::ASCENDING}, {}, stream); + test_sort_two_tables(preprocessed_lhs, preprocessed_empty_rhs); + } +} + template struct NaNTableViewTest : public cudf::test::BaseFixture {};