diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 9750ed7c5c4..db560077298 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -470,6 +470,7 @@ add_library( src/structs/structs_column_factories.cu src/structs/structs_column_view.cpp src/structs/utilities.cpp + src/table/row_operators.cu src/table/table.cpp src/table/table_device_view.cu src/table/table_view.cpp diff --git a/cpp/benchmarks/CMakeLists.txt b/cpp/benchmarks/CMakeLists.txt index 9e8a632a3ae..d863e6e05a9 100644 --- a/cpp/benchmarks/CMakeLists.txt +++ b/cpp/benchmarks/CMakeLists.txt @@ -164,6 +164,7 @@ ConfigureBench(SEARCH_BENCH search/search.cpp) # ################################################################################################## # * sort benchmark -------------------------------------------------------------------------------- ConfigureBench(SORT_BENCH sort/rank.cpp sort/sort.cpp sort/sort_strings.cpp) +ConfigureNVBench(SORT_NVBENCH sort/sort_structs.cpp) # ################################################################################################## # * quantiles benchmark diff --git a/cpp/benchmarks/sort/sort_structs.cpp b/cpp/benchmarks/sort/sort_structs.cpp new file mode 100644 index 00000000000..d2d8e12c377 --- /dev/null +++ b/cpp/benchmarks/sort/sort_structs.cpp @@ -0,0 +1,84 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include + +#include + +#include +#include + +#include + +#include + +void nvbench_sort_struct(nvbench::state& state) +{ + cudf::rmm_pool_raii pool_raii; + + using Type = int; + using column_wrapper = cudf::test::fixed_width_column_wrapper; + std::default_random_engine generator; + std::uniform_int_distribution distribution(0, 100); + + const cudf::size_type n_rows{static_cast(state.get_int64("NumRows"))}; + const cudf::size_type n_cols{1}; + const cudf::size_type depth{static_cast(state.get_int64("Depth"))}; + const bool nulls{static_cast(state.get_int64("Nulls"))}; + + // Create columns with values in the range [0,100) + std::vector columns; + columns.reserve(n_cols); + std::generate_n(std::back_inserter(columns), n_cols, [&]() { + auto const elements = cudf::detail::make_counting_transform_iterator( + 0, [&](auto row) { return distribution(generator); }); + if (!nulls) return column_wrapper(elements, elements + n_rows); + auto valids = + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i % 10 != 0; }); + return column_wrapper(elements, elements + n_rows, valids); + }); + + std::vector> cols; + std::transform(columns.begin(), columns.end(), std::back_inserter(cols), [](column_wrapper& col) { + return col.release(); + }); + + std::vector> child_cols = std::move(cols); + // Lets add some layers + for (int i = 0; i < depth; i++) { + std::vector struct_validity; + std::uniform_int_distribution bool_distribution(0, 100 * (i + 1)); + std::generate_n( + std::back_inserter(struct_validity), n_rows, [&]() { return bool_distribution(generator); }); + cudf::test::structs_column_wrapper struct_col(std::move(child_cols), struct_validity); + child_cols = std::vector>{}; + child_cols.push_back(struct_col.release()); + } + + // Create table view + auto const input = cudf::table(std::move(child_cols)); + + state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) { + rmm::cuda_stream_view stream_view{launch.get_stream()}; + cudf::detail::sorted_order(input, {}, {}, stream_view, rmm::mr::get_current_device_resource()); + }); +} + +NVBENCH_BENCH(nvbench_sort_struct) + .set_name("sort_struct") + .add_int64_power_of_two_axis("NumRows", {10, 18, 26}) + .add_int64_axis("Depth", {1, 8}) + .add_int64_axis("Nulls", {0, 1}); diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh new file mode 100644 index 00000000000..d305cfa66b5 --- /dev/null +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -0,0 +1,424 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * 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 +#include + +#include +#include +#include + +#include +#include + +#include +#include +#include + +namespace cudf { +namespace experimental { + +/** + * @brief A map from cudf::type_id to cudf type that excludes LIST and STRUCT types. + * + * To be used with type_dispatcher in place of the default map, when it is required that STRUCT and + * LIST map to void. This is useful when we want to avoid recursion in a functor. For example, in + * element_comparator, we have a specialization for STRUCT but the type_dispatcher in it is only + * used to dispatch to the same functor for non-nested types. Even when we're guaranteed to not have + * non-nested types at that point, the compiler doesn't know this and would try to create recursive + * code which is very slow. + * + * Usage: + * @code + * type_dispatcher(data_type(), functor{}); + * @endcode + */ +template +struct dispatch_void_if_nested { + using type = std::conditional_t>; +}; + +namespace row { + +namespace lexicographic { + +/** + * @brief Computes whether one row is lexicographically *less* than another row. + * + * Lexicographic ordering is determined by: + * - Two rows are compared element by element. + * - The first mismatching element defines which row is lexicographically less + * or greater than the other. + * + * Lexicographic ordering is exactly equivalent to doing an alphabetical sort of + * two words, for example, `aac` would be *less* than (or precede) `abb`. The + * second letter in both words is the first non-equal letter, and `a < b`, thus + * `aac < abb`. + * + * @tparam Nullate A cudf::nullate type describing how to check for nulls. + */ +template +class device_row_comparator { + friend class self_comparator; + + /** + * @brief Construct a function object for performing a lexicographic + * comparison between the rows of two tables. + * + * @param has_nulls Indicates if either input table contains columns with nulls. + * @param lhs The first table + * @param rhs The second table (may be the same table as `lhs`) + * @param depth Optional, device array the same length as a row that contains starting depths of + * columns if they're nested, and 0 otherwise. + * @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 `nullopt`, 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 `nullopt`, then null precedence would be + * `null_order::BEFORE` for all columns. + */ + device_row_comparator( + Nullate has_nulls, + table_device_view lhs, + table_device_view rhs, + std::optional> depth = std::nullopt, + std::optional> column_order = std::nullopt, + std::optional> null_precedence = std::nullopt) noexcept + : _lhs{lhs}, + _rhs{rhs}, + _nulls{has_nulls}, + _depth{depth}, + _column_order{column_order}, + _null_precedence{null_precedence} + { + } + + /** + * @brief Performs a relational comparison between two elements in two columns. + */ + class element_comparator { + public: + /** + * @brief Construct type-dispatched function object for performing a + * relational comparison between two elements. + * + * @note `lhs` and `rhs` may be the same. + * + * @param has_nulls Indicates if either input column contains nulls. + * @param lhs The column containing the first element + * @param rhs The column containing the second element (may be the same as lhs) + * @param null_precedence Indicates how null values are ordered with other values + * @param depth The depth of the column if part of a nested column @see + * preprocessed_table::depths + */ + __device__ element_comparator(Nullate has_nulls, + column_device_view lhs, + column_device_view rhs, + null_order null_precedence = null_order::BEFORE, + int depth = 0) + : _lhs{lhs}, _rhs{rhs}, _nulls{has_nulls}, _null_precedence{null_precedence}, _depth{depth} + { + } + + /** + * @brief Performs a relational comparison between the specified elements + * + * @param lhs_element_index The index of the first element + * @param rhs_element_index The index of the second element + * @return Indicates the relationship between the elements in the `lhs` and `rhs` columns, along + * with the depth at which a null value was encountered. + */ + template ())> + __device__ cuda::std::pair operator()( + size_type const lhs_element_index, size_type const rhs_element_index) const noexcept + { + if (_nulls) { + bool const lhs_is_null{_lhs.is_null(lhs_element_index)}; + bool const rhs_is_null{_rhs.is_null(rhs_element_index)}; + + if (lhs_is_null or rhs_is_null) { // at least one is null + return cuda::std::make_pair(null_compare(lhs_is_null, rhs_is_null, _null_precedence), + _depth); + } + } + + return cuda::std::make_pair(relational_compare(_lhs.element(lhs_element_index), + _rhs.element(rhs_element_index)), + std::numeric_limits::max()); + } + + template () and + not std::is_same_v)> + __device__ cuda::std::pair operator()(size_type const lhs_element_index, + size_type const rhs_element_index) + { + // TODO: make this CUDF_UNREACHABLE + cudf_assert(false && "Attempted to compare elements of uncomparable types."); + return cuda::std::make_pair(weak_ordering::LESS, std::numeric_limits::max()); + } + + template )> + __device__ cuda::std::pair operator()(size_type const lhs_element_index, + size_type const rhs_element_index) + { + column_device_view lcol = _lhs; + column_device_view rcol = _rhs; + int depth = _depth; + while (lcol.type().id() == type_id::STRUCT) { + bool const lhs_is_null{lcol.is_null(lhs_element_index)}; + bool const rhs_is_null{rcol.is_null(rhs_element_index)}; + + if (lhs_is_null or rhs_is_null) { // at least one is null + weak_ordering state = null_compare(lhs_is_null, rhs_is_null, _null_precedence); + return cuda::std::make_pair(state, depth); + } + + // Structs have been modified to only have 1 child when using this. + lcol = lcol.children()[0]; + rcol = rcol.children()[0]; + ++depth; + } + + auto const comparator = element_comparator{_nulls, lcol, rcol, _null_precedence, depth}; + return cudf::type_dispatcher( + lcol.type(), comparator, lhs_element_index, rhs_element_index); + } + + private: + column_device_view const _lhs; + column_device_view const _rhs; + Nullate const _nulls; + null_order const _null_precedence; + int const _depth; + }; + + public: + /** + * @brief Checks whether the row at `lhs_index` in the `lhs` table compares + * lexicographically less than the row at `rhs_index` in the `rhs` table. + * + * @param lhs_index The index of row in the `lhs` table to examine + * @param rhs_index The index of the row in the `rhs` table to examine + * @return `true` if row from the `lhs` table compares less than row in the `rhs` table + */ + __device__ bool operator()(size_type const lhs_index, size_type const rhs_index) const noexcept + { + int last_null_depth = std::numeric_limits::max(); + for (size_type i = 0; i < _lhs.num_columns(); ++i) { + int const depth = _depth.has_value() ? (*_depth)[i] : 0; + if (depth > last_null_depth) { continue; } + + bool const ascending = + _column_order.has_value() ? (*_column_order)[i] == order::ASCENDING : true; + + null_order const null_precedence = + _null_precedence.has_value() ? (*_null_precedence)[i] : null_order::BEFORE; + + auto const comparator = + element_comparator{_nulls, _lhs.column(i), _rhs.column(i), null_precedence, depth}; + + weak_ordering state; + cuda::std::tie(state, last_null_depth) = + cudf::type_dispatcher(_lhs.column(i).type(), comparator, lhs_index, rhs_index); + + if (state == weak_ordering::EQUIVALENT) { continue; } + + return state == (ascending ? weak_ordering::LESS : weak_ordering::GREATER); + } + return false; + } + + private: + table_device_view const _lhs; + table_device_view const _rhs; + Nullate const _nulls{}; + std::optional> const _depth; + std::optional> const _column_order; + std::optional> const _null_precedence; +}; // class device_row_comparator + +struct preprocessed_table { + using table_device_view_owner = + std::invoke_result_t; + + /** + * @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 `lex::self_comparator` to avoid preprocessing again. + * + * @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. + */ + static std::shared_ptr create(table_view const& table, + host_span column_order, + host_span null_precedence, + rmm::cuda_stream_view stream); + + private: + friend class self_comparator; + + 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)){}; + + /** + * @brief Implicit conversion operator to a `table_device_view` of the preprocessed table. + * + * @return table_device_view + */ + operator table_device_view() { return *_t; } + + /** + * @brief Get a device array containing the desired order of each column in the preprocessed table + * + * @return std::optional> Device array containing respective column + * orders. If no explicit column orders were specified during the creation of this object then + * this will be `nullopt`. + */ + [[nodiscard]] std::optional> column_order() const + { + return _column_order.size() ? std::optional>(_column_order) + : std::nullopt; + } + + /** + * @brief Get a device array containing the desired null precedence of each column in the + * preprocessed table + * + * @return std::optional> Device array containing respective column + * null precedence. If no explicit column null precedences were specified during the creation of + * this object then this will be `nullopt`. + */ + [[nodiscard]] std::optional> null_precedence() const + { + return _null_precedence.size() ? std::optional>(_null_precedence) + : std::nullopt; + } + + /** + * @brief Get a device array containing the depth of each column in the preprocessed table + * + * @see struct_linearize() + * + * @return std::optional> Device array containing respective column depths. + * If there are no nested columns in the table then this will be `nullopt`. + */ + [[nodiscard]] std::optional> depths() const + { + return _depths.size() ? std::optional>(_depths) : std::nullopt; + } + + private: + table_device_view_owner _t; + rmm::device_uvector _column_order; + rmm::device_uvector _null_precedence; + rmm::device_uvector _depths; +}; + +/** + * @brief An owning object that can be used to lexicographically compare two rows of the same table + * + * This class can take a table_view and preprocess certain columns to allow for lexicographical + * comparison. The preprocessed table and temporary data required for the comparison are created and + * owned by this class. + * + * Alternatively, `self_comparator` can be constructed from an existing + * `shared_ptr` when sharing the same table among multiple comparators. + * + * This class can then provide a functor object that can used on the device. + * The object of this class must outlive the usage of the device functor. + */ +class self_comparator { + public: + /** + * @brief Construct an owning object for performing a lexicographic comparison between two rows of + * the same table. + * + * @param table 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. + * @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. + * @param stream The stream to construct this object on. Not the stream that will be used for + * comparisons using this object. + */ + self_comparator(table_view const& t, + host_span column_order = {}, + host_span null_precedence = {}, + rmm::cuda_stream_view stream = rmm::cuda_stream_default) + : d_t{preprocessed_table::create(t, column_order, null_precedence, stream)} + { + } + + /** + * @brief Construct an owning object for performing a lexicographic comparison between two rows of + * the same preprocessed table. + * + * This constructor allows independently constructing a `preprocessed_table` and sharing it among + * multiple comparators. + * + * @param t A table preprocessed for lexicographic comparison + */ + self_comparator(std::shared_ptr t) : d_t{std::move(t)} {} + + /** + * @brief Return the binary operator for comparing rows in the table. + * + * Returns a binary callable, `F`, with signature `bool F(size_t, size_t)`. + * + * `F(i,j)` returns true if and only if row `i` compares lexicographically less than row `j`. + * + * @tparam Nullate Optional, A cudf::nullate type describing how to check for nulls. + */ + template + device_row_comparator device_comparator(Nullate nullate = {}) const + { + return device_row_comparator( + nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence()); + } + + private: + std::shared_ptr d_t; +}; + +} // namespace lexicographic +} // namespace row +} // namespace experimental +} // namespace cudf diff --git a/cpp/include/cudf/utilities/traits.hpp b/cpp/include/cudf/utilities/traits.hpp index 504ec6de405..ed24517f55b 100644 --- a/cpp/include/cudf/utilities/traits.hpp +++ b/cpp/include/cudf/utilities/traits.hpp @@ -124,6 +124,31 @@ constexpr inline bool is_relationally_comparable() return detail::is_relationally_comparable_impl::value; } +namespace detail { +/** + * @brief Helper functor to check if a specified type `T` supports relational comparisons. + */ +struct unary_relationally_comparable_functor { + template + inline bool operator()() const + { + return cudf::is_relationally_comparable(); + } +}; +} // namespace detail + +/** + * @brief Checks whether `data_type` `type` supports relational comparisons. + * + * @param type Data_type for comparison. + * @return true If `type` supports relational comparisons. + * @return false If `type` does not support relational comparisons. + */ +inline bool is_relationally_comparable(data_type type) +{ + return type_dispatcher(type, detail::unary_relationally_comparable_functor{}); +} + /** * @brief Indicates whether objects of types `L` and `R` can be compared * for equality. diff --git a/cpp/src/sort/sort_impl.cuh b/cpp/src/sort/sort_impl.cuh index 881503a49e3..2f093fd7d2d 100644 --- a/cpp/src/sort/sort_impl.cuh +++ b/cpp/src/sort/sort_impl.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -20,6 +20,7 @@ #include #include #include +#include #include #include #include @@ -123,16 +124,10 @@ std::unique_ptr sorted_order(table_view input, mutable_indices_view.end(), 0); - auto flattened = structs::detail::flatten_nested_columns(input, column_order, null_precedence); - auto device_table = table_device_view::create(flattened, stream); - auto const d_column_order = make_device_uvector_async(flattened.orders(), stream); + auto comp = + experimental::row::lexicographic::self_comparator(input, column_order, null_precedence, stream); + auto comparator = comp.device_comparator(nullate::DYNAMIC{has_nested_nulls(input)}); - auto const d_null_precedence = make_device_uvector_async(flattened.null_orders(), stream); - auto const comparator = row_lexicographic_comparator(nullate::DYNAMIC{has_nulls(flattened)}, - *device_table, - *device_table, - d_column_order.data(), - d_null_precedence.data()); if (stable) { thrust::stable_sort(rmm::exec_policy(stream), mutable_indices_view.begin(), diff --git a/cpp/src/table/row_operators.cu b/cpp/src/table/row_operators.cu new file mode 100644 index 00000000000..0a9396ccdf7 --- /dev/null +++ b/cpp/src/table/row_operators.cu @@ -0,0 +1,207 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include +#include +#include +#include + +#include + +namespace cudf { +namespace experimental { + +namespace { + +/** + * @brief Decompose all struct columns in a table + * + * If a struct column is a tree with N leaves, then this function decomposes the tree into + * N "linear trees" (branch factor == 1) and prunes common parents. Also returns a vector of + * per-column `depth`s. + * + * A `depth` value is the number of nested levels as parent of the column in the original, + * non-decomposed table, which are pruned during decomposition. + * + * For example, if the original table has a column `Struct, decimal>`, + * S1 + * / \ + * S2 d + * / \ + * i f + * then after decomposition, we get three columns: + * `Struct>`, `float`, and `decimal`. + * 0 2 1 <- depths + * S1 + * | + * S2 d + * | + * i f + * The depth of the first column is 0 because it contains all its parent levels, while the depth + * of the second column is 2 because two of its parent struct levels were pruned. + * + * Similarly, a struct column of type Struct> is decomposed as follows + * S1 + * / \ + * i S2 + * / \ + * f d + * + * 0 1 2 <- depths + * S1 S2 d + * | | + * i f + * + * @param table The table whose struct columns to decompose. + * @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, + host_span column_order = {}, + host_span null_precedence = {}) +{ + std::vector verticalized_columns; + std::vector new_column_order; + std::vector new_null_precedence; + std::vector verticalized_col_depths; + for (size_type col_idx = 0; col_idx < table.num_columns(); ++col_idx) { + auto const& col = table.column(col_idx); + if (is_nested(col.type())) { + // convert and insert + std::vector r_verticalized_columns; + std::vector r_verticalized_col_depths; + std::vector flattened; + std::vector depths; + // TODO: Here I added a bogus leaf column at the beginning to help in the while loop below. + // Refactor the while loop so that it can handle the last case. + flattened.push_back(make_empty_column(type_id::INT32)->view()); + std::function recursive_child = [&](column_view const& c, + int depth) { + flattened.push_back(c); + depths.push_back(depth); + if (c.type().id() == type_id::STRUCT) { + for (int child_idx = 0; child_idx < c.num_children(); ++child_idx) { + auto scol = structs_column_view(c); + recursive_child(scol.get_sliced_child(child_idx), depth + 1); + } + } + }; + recursive_child(col, 0); + int curr_col_idx = flattened.size() - 1; + column_view curr_col = flattened[curr_col_idx]; + while (curr_col_idx > 0) { + auto const& prev_col = flattened[curr_col_idx - 1]; + if (not is_nested(prev_col.type())) { + // We hit a column that's a leaf so seal this hierarchy + r_verticalized_columns.push_back(curr_col); + r_verticalized_col_depths.push_back(depths[curr_col_idx - 1]); + curr_col = prev_col; + } else { + curr_col = column_view(prev_col.type(), + prev_col.size(), + nullptr, + prev_col.null_mask(), + UNKNOWN_NULL_COUNT, + prev_col.offset(), + {curr_col}); + } + --curr_col_idx; + } + verticalized_columns.insert( + verticalized_columns.end(), r_verticalized_columns.rbegin(), r_verticalized_columns.rend()); + verticalized_col_depths.insert(verticalized_col_depths.end(), + r_verticalized_col_depths.rbegin(), + r_verticalized_col_depths.rend()); + if (not column_order.empty()) { + new_column_order.insert( + new_column_order.end(), r_verticalized_columns.size(), column_order[col_idx]); + } + if (not null_precedence.empty()) { + new_null_precedence.insert( + new_null_precedence.end(), r_verticalized_columns.size(), null_precedence[col_idx]); + } + } else { + verticalized_columns.push_back(col); + verticalized_col_depths.push_back(0); + if (not column_order.empty()) { new_column_order.push_back(column_order[col_idx]); } + if (not null_precedence.empty()) { new_null_precedence.push_back(null_precedence[col_idx]); } + } + } + return std::make_tuple(table_view(verticalized_columns), + std::move(new_column_order), + std::move(new_null_precedence), + std::move(verticalized_col_depths)); +} + +/** + * @brief Check a table for compatibility with lexicographic comparison + * + * Checks whether a given table contains columns of non-relationally comparable types. + */ +void check_lex_compatibility(table_view const& input) +{ + // Basically check if there's any LIST hiding anywhere in the table + std::function check_column = [&](column_view const& c) { + CUDF_EXPECTS(c.type().id() != type_id::LIST, + "Cannot lexicographic compare a table with a LIST column"); + if (not is_nested(c.type())) { + CUDF_EXPECTS(is_relationally_comparable(c.type()), + "Cannot lexicographic compare a table with a column of type " + + jit::get_type_name(c.type())); + } + for (auto child = c.child_begin(); child < c.child_end(); ++child) { + check_column(*child); + } + }; + for (column_view const& c : input) { + check_column(c); + } +} + +} // namespace + +namespace row { + +namespace lexicographic { + +std::shared_ptr preprocessed_table::create( + table_view const& t, + host_span column_order, + host_span null_precedence, + 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); + + auto d_t = table_device_view::create(verticalized_lhs, stream); + auto d_column_order = detail::make_device_uvector_async(new_column_order, stream); + auto d_null_precedence = detail::make_device_uvector_async(new_null_precedence, stream); + auto d_depths = detail::make_device_uvector_async(verticalized_col_depths, stream); + + 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))); +} + +} // namespace lexicographic +} // namespace row +} // namespace experimental +} // namespace cudf diff --git a/cpp/src/table/table_view.cpp b/cpp/src/table/table_view.cpp index 365ff67263c..c89906f3480 100644 --- a/cpp/src/table/table_view.cpp +++ b/cpp/src/table/table_view.cpp @@ -97,15 +97,6 @@ table_view scatter_columns(table_view const& source, } namespace detail { -namespace { -struct is_relationally_comparable_functor { - template - constexpr bool operator()() - { - return cudf::is_relationally_comparable(); - } -}; -} // namespace template bool is_relationally_comparable(TableView const& lhs, TableView const& rhs) @@ -114,8 +105,7 @@ bool is_relationally_comparable(TableView const& lhs, TableView const& rhs) thrust::counting_iterator(lhs.num_columns()), [lhs, rhs](auto const i) { return lhs.column(i).type() == rhs.column(i).type() and - type_dispatcher(lhs.column(i).type(), - is_relationally_comparable_functor{}); + cudf::is_relationally_comparable(lhs.column(i).type()); }); } diff --git a/cpp/tests/sort/sort_test.cpp b/cpp/tests/sort/sort_test.cpp index 7f9f40e98b8..ed86277cd2b 100644 --- a/cpp/tests/sort/sort_test.cpp +++ b/cpp/tests/sort/sort_test.cpp @@ -290,6 +290,100 @@ TYPED_TEST(Sort, WithNestedStructColumn) } } +TYPED_TEST(Sort, WithNullableStructColumn) +{ + // Test for a struct column that has nulls on struct layer but not pushed down on the child + using T = int; + using fwcw = cudf::test::fixed_width_column_wrapper; + using mask = std::vector; + + auto make_struct = [&](std::vector> child_cols, + std::vector nulls) { + cudf::test::structs_column_wrapper struct_col(std::move(child_cols)); + auto struct_ = struct_col.release(); + struct_->set_null_mask(cudf::test::detail::make_null_mask(nulls.begin(), nulls.end())); + return struct_; + }; + + { + /* + /+-------------+ + |s1{s2{a,b}, c}| + +--------------+ + 0 | { {1, 1}, 5}| + 1 | { {1, 2}, 4}| + 2 | {@{2, 1}, 6}| + 3 | {@{2, 2}, 5}| + 4 | @{ {2, 2}, 3}| + 5 | @{ {1, 1}, 3}| + 6 | { {1, 2}, 3}| + 7 | {@{1, 1}, 4}| + 8 | { {2, 1}, 5}| + +--------------+ + + Intermediate representation: + s1{s2{a}}, b, c + */ + + auto col_a = fwcw{1, 1, 2, 2, 2, 1, 1, 1, 2}; + auto col_b = fwcw{1, 2, 1, 2, 2, 1, 2, 1, 1}; + auto s2_mask = mask{1, 1, 0, 0, 1, 1, 1, 0, 1}; + auto col_c = fwcw{5, 4, 6, 5, 3, 3, 3, 4, 5}; + auto s1_mask = mask{1, 1, 1, 1, 0, 0, 1, 1, 1}; + + std::vector> s2_children; + s2_children.push_back(col_a.release()); + s2_children.push_back(col_b.release()); + auto s2 = make_struct(std::move(s2_children), s2_mask); + + std::vector> s1_children; + s1_children.push_back(std::move(s2)); + s1_children.push_back(col_c.release()); + auto s1 = make_struct(std::move(s1_children), s1_mask); + + auto expect = fwcw{4, 5, 7, 3, 2, 0, 6, 1, 8}; + run_sort_test(table_view({s1->view()}), expect); + } + { /* + /+-------------+ + |s1{a,s2{b, c}}| + +--------------+ + 0 | {1, {1, 5}}| + 1 | {1, {2, 4}}| + 2 | {2, @{2, 6}}| + 3 | {2, @{1, 5}}| + 4 | @{2, {2, 3}}| + 5 | @{1, {1, 3}}| + 6 | {1, {2, 3}}| + 7 | {1, @{1, 4}}| + 8 | {2, {1, 5}}| + +--------------+ + + Intermediate representation: + s1{a}, s2{b}, c + */ + + auto s1_mask = mask{1, 1, 1, 1, 0, 0, 1, 1, 1}; + auto col_a = fwcw{1, 1, 2, 2, 2, 1, 1, 1, 2}; + auto s2_mask = mask{1, 1, 0, 0, 1, 1, 1, 0, 1}; + auto col_b = fwcw{1, 2, 1, 2, 2, 1, 2, 1, 1}; + auto col_c = fwcw{5, 4, 6, 5, 3, 3, 3, 4, 5}; + + std::vector> s22_children; + s22_children.push_back(col_b.release()); + s22_children.push_back(col_c.release()); + auto s22 = make_struct(std::move(s22_children), s2_mask); + + std::vector> s12_children; + s12_children.push_back(col_a.release()); + s12_children.push_back(std::move(s22)); + auto s12 = make_struct(std::move(s12_children), s1_mask); + + auto expect = fwcw{4, 5, 7, 0, 6, 1, 2, 3, 8}; + run_sort_test(table_view({s12->view()}), expect); + } +} + TYPED_TEST(Sort, WithSingleStructColumn) { using T = TypeParam; @@ -641,6 +735,20 @@ TYPED_TEST(Sort, ZeroSizedColumns) run_sort_test(input, expected, column_order); } +TYPED_TEST(Sort, WithListColumn) +{ + using T = int; + lists_column_wrapper lc{{1, 2, 3}, {4, 5, 6}, {7, 8, 9}}; + CUDF_EXPECT_THROW_MESSAGE(cudf::sort(table_view({lc})), + "Cannot lexicographic compare a table with a LIST column"); + + std::vector> child_cols; + child_cols.push_back(lc.release()); + structs_column_wrapper sc{std::move(child_cols), {1, 0, 1}}; + CUDF_EXPECT_THROW_MESSAGE(cudf::sort(table_view({sc})), + "Cannot lexicographic compare a table with a LIST column"); +} + struct SortByKey : public BaseFixture { };