From 399d5b5d996db48074804c079b59fb529356cfd3 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Mon, 18 Oct 2021 15:39:23 -0700 Subject: [PATCH 1/4] Small clean up to simplify column selection code in ORC reader (#9444) Remove `has_nested_column` from column selection and the `impl` class. Deduce the existence of nested columns from the depth of the schema instead. Authors: - Vukasin Milovanovic (https://github.com/vuule) Approvers: - Nghia Truong (https://github.com/ttnghia) - Vyas Ramasubramani (https://github.com/vyasr) URL: https://github.com/rapidsai/cudf/pull/9444 --- cpp/src/io/orc/reader_impl.cu | 34 +++++++++++++--------------------- cpp/src/io/orc/reader_impl.hpp | 1 - 2 files changed, 13 insertions(+), 22 deletions(-) diff --git a/cpp/src/io/orc/reader_impl.cu b/cpp/src/io/orc/reader_impl.cu index 00ae00a9171..d1c8c3661f4 100644 --- a/cpp/src/io/orc/reader_impl.cu +++ b/cpp/src/io/orc/reader_impl.cu @@ -467,29 +467,21 @@ class aggregate_orc_metadata { * @param level current level of nesting. * @param id current column id that needs to be added. * @param has_timestamp_column True if timestamp column present and false otherwise. - * @param has_nested_column True if any of the selected column is a nested type. */ void add_column(std::vector>& selection, std::vector const& types, const size_t level, const uint32_t id, - bool& has_timestamp_column, - bool& has_nested_column) + bool& has_timestamp_column) { if (level == selection.size()) { selection.emplace_back(); } - selection[level].push_back({id, 0}); - const int col_id = selection[level].size() - 1; + selection[level].push_back({id, static_cast(types[id].subtypes.size())}); if (types[id].kind == orc::TIMESTAMP) { has_timestamp_column = true; } - if (types[id].kind == orc::MAP or types[id].kind == orc::LIST or - types[id].kind == orc::STRUCT) { - has_nested_column = true; - for (const auto child_id : types[id].subtypes) { - // Since nested column needs to be processed before its child can be processed, - // child column is being added to next level - add_column(selection, types, level + 1, child_id, has_timestamp_column, has_nested_column); - } - selection[level][col_id].num_children = types[id].subtypes.size(); + for (const auto child_id : types[id].subtypes) { + // Since nested column needs to be processed before its child can be processed, + // child column is being added to next level + add_column(selection, types, level + 1, child_id, has_timestamp_column); } } @@ -498,12 +490,11 @@ class aggregate_orc_metadata { * * @param use_names List of column names to select * @param has_timestamp_column True if timestamp column present and false otherwise - * @param has_nested_column True if any of the selected column is a nested type. * * @return Vector of list of ORC column meta-data */ std::vector> select_columns( - std::vector const& use_names, bool& has_timestamp_column, bool& has_nested_column) + std::vector const& use_names, bool& has_timestamp_column) { auto const& pfm = per_file_metadata[0]; std::vector> selection; @@ -520,7 +511,7 @@ class aggregate_orc_metadata { auto col_id = pfm.ff.types[0].subtypes[index]; if (pfm.get_column_name(col_id) == use_name) { name_found = true; - add_column(selection, pfm.ff.types, 0, col_id, has_timestamp_column, has_nested_column); + add_column(selection, pfm.ff.types, 0, col_id, has_timestamp_column); // Should start with next index index = i + 1; break; @@ -530,7 +521,7 @@ class aggregate_orc_metadata { } } else { for (auto const& col_id : pfm.ff.types[0].subtypes) { - add_column(selection, pfm.ff.types, 0, col_id, has_timestamp_column, has_nested_column); + add_column(selection, pfm.ff.types, 0, col_id, has_timestamp_column); } } @@ -1164,8 +1155,7 @@ reader::impl::impl(std::vector>&& sources, _metadata = std::make_unique(_sources); // Select only columns required by the options - _selected_columns = - _metadata->select_columns(options.get_columns(), _has_timestamp_column, _has_nested_column); + _selected_columns = _metadata->select_columns(options.get_columns(), _has_timestamp_column); // Override output timestamp resolution if requested if (options.get_timestamp_type().id() != type_id::EMPTY) { @@ -1187,7 +1177,9 @@ table_with_metadata reader::impl::read(size_type skip_rows, const std::vector>& stripes, rmm::cuda_stream_view stream) { - CUDF_EXPECTS(skip_rows == 0 or (not _has_nested_column), + // Selected columns at different levels of nesting are stored in different elements + // of `_selected_columns`; thus, size == 1 means no nested columns + CUDF_EXPECTS(skip_rows == 0 or _selected_columns.size() == 1, "skip_rows is not supported by nested columns"); std::vector> out_columns; diff --git a/cpp/src/io/orc/reader_impl.hpp b/cpp/src/io/orc/reader_impl.hpp index 4be75b6cc2a..8fed64ed64c 100644 --- a/cpp/src/io/orc/reader_impl.hpp +++ b/cpp/src/io/orc/reader_impl.hpp @@ -215,7 +215,6 @@ class reader::impl { bool _use_index = true; bool _use_np_dtypes = true; bool _has_timestamp_column = false; - bool _has_nested_column = false; std::vector _decimal_cols_as_float; data_type _timestamp_type{type_id::EMPTY}; reader_column_meta _col_meta; From 823958b0e688ea366235be7662677238af1fb8ed Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Mon, 18 Oct 2021 16:41:24 -0600 Subject: [PATCH 2/4] Implement `lists::stable_sort_lists` for stable sorting of elements within each row of lists column (#9425) This PR adds `lists::stable_sort_lists` that can sort elements within rows of lists column using stable sort. This is necessary for implementing `lists::drop_list_duplicates` that operates on keys-values columns input when we want to remove the values corresponding to duplicate keys with `KEEP_FIRST` or `KEEP_LAST` option. In order to implement `lists::stable_sort_lists`, stable sort versions for the `segmented_sorted_order` and `segmented_sort_by_key` have also been implemented, which can maintain the order of equally-compared elements within segments. This PR blocks #9345. Authors: - Nghia Truong (https://github.com/ttnghia) Approvers: - Conor Hoekstra (https://github.com/codereport) - MithunR (https://github.com/mythrocks) URL: https://github.com/rapidsai/cudf/pull/9425 --- cpp/include/cudf/detail/sorting.hpp | 27 +++ cpp/include/cudf/lists/detail/sorting.hpp | 13 ++ cpp/include/cudf/lists/sorting.hpp | 12 ++ cpp/include/cudf/sorting.hpp | 25 +++ cpp/src/lists/segmented_sort.cu | 44 +++++ cpp/src/sort/segmented_sort.cu | 168 ++++++++++++++---- cpp/tests/lists/sort_lists_tests.cpp | 205 ++++++++++++++-------- 7 files changed, 389 insertions(+), 105 deletions(-) diff --git a/cpp/include/cudf/detail/sorting.hpp b/cpp/include/cudf/detail/sorting.hpp index b5dfb34c043..3aa85e87b1d 100644 --- a/cpp/include/cudf/detail/sorting.hpp +++ b/cpp/include/cudf/detail/sorting.hpp @@ -76,6 +76,19 @@ std::unique_ptr segmented_sorted_order( rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +/** + * @copydoc cudf::stable_segmented_sorted_order + * + * @param[in] stream CUDA stream used for device memory operations and kernel launches. + */ +std::unique_ptr stable_segmented_sorted_order( + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order = {}, + std::vector const& null_precedence = {}, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + /** * @copydoc cudf::segmented_sort_by_key * @@ -90,6 +103,20 @@ std::unique_ptr segmented_sort_by_key( rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +/** + * @copydoc cudf::stable_segmented_sort_by_key + * + * @param[in] stream CUDA stream used for device memory operations and kernel launches. + */ +std::unique_ptr
stable_segmented_sort_by_key( + table_view const& values, + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order = {}, + std::vector const& null_precedence = {}, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + /** * @copydoc cudf::sort * diff --git a/cpp/include/cudf/lists/detail/sorting.hpp b/cpp/include/cudf/lists/detail/sorting.hpp index f68ff872020..1068a4c4b69 100644 --- a/cpp/include/cudf/lists/detail/sorting.hpp +++ b/cpp/include/cudf/lists/detail/sorting.hpp @@ -34,6 +34,19 @@ std::unique_ptr sort_lists( null_order null_precedence, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +/** + * @copydoc cudf::lists::stable_sort_lists + * + * @param stream CUDA stream used for device memory operations and kernel launches. + */ +std::unique_ptr stable_sort_lists( + lists_column_view const& input, + order column_order, + null_order null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + } // namespace detail } // namespace lists } // namespace cudf diff --git a/cpp/include/cudf/lists/sorting.hpp b/cpp/include/cudf/lists/sorting.hpp index e27f3d03d86..55fd722ca14 100644 --- a/cpp/include/cudf/lists/sorting.hpp +++ b/cpp/include/cudf/lists/sorting.hpp @@ -54,6 +54,18 @@ std::unique_ptr sort_lists( null_order null_precedence, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +/** + * @brief Segmented sort of the elements within a list in each row of a list column using stable + * sort. + * + * @copydoc cudf::lists::sort_lists + */ +std::unique_ptr stable_sort_lists( + lists_column_view const& source_column, + order column_order, + null_order null_precedence, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + /** @} */ // end of group } // namespace lists } // namespace cudf diff --git a/cpp/include/cudf/sorting.hpp b/cpp/include/cudf/sorting.hpp index 2c9c9c64a64..c17abe8267d 100644 --- a/cpp/include/cudf/sorting.hpp +++ b/cpp/include/cudf/sorting.hpp @@ -211,6 +211,18 @@ std::unique_ptr segmented_sorted_order( std::vector const& null_precedence = {}, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +/** + * @brief Returns sorted order after stably sorting each segment in the table. + * + * @copydoc cudf::segmented_sorted_order + */ +std::unique_ptr stable_segmented_sorted_order( + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order = {}, + std::vector const& null_precedence = {}, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + /** * @brief Performs a lexicographic segmented sort of a table * @@ -241,5 +253,18 @@ std::unique_ptr
segmented_sort_by_key( std::vector const& null_precedence = {}, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +/** + * @brief Performs a stably lexicographic segmented sort of a table + * + * @copydoc cudf::segmented_sort_by_key + */ +std::unique_ptr
stable_segmented_sort_by_key( + table_view const& values, + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order = {}, + std::vector const& null_precedence = {}, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + /** @} */ // end of group } // namespace cudf diff --git a/cpp/src/lists/segmented_sort.cu b/cpp/src/lists/segmented_sort.cu index b085d1e77d1..088db226c24 100644 --- a/cpp/src/lists/segmented_sort.cu +++ b/cpp/src/lists/segmented_sort.cu @@ -268,6 +268,40 @@ std::unique_ptr sort_lists(lists_column_view const& input, input.null_count(), std::move(null_mask)); } + +std::unique_ptr stable_sort_lists(lists_column_view const& input, + order column_order, + null_order null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + if (input.is_empty()) { return empty_like(input.parent()); } + + auto output_offset = make_numeric_column( + input.offsets().type(), input.size() + 1, mask_state::UNALLOCATED, stream, mr); + thrust::transform(rmm::exec_policy(stream), + input.offsets_begin(), + input.offsets_end(), + output_offset->mutable_view().template begin(), + [first = input.offsets_begin()] __device__(auto offset_index) { + return offset_index - *first; + }); + + auto const child = input.get_sliced_child(stream); + auto const sorted_child_table = stable_segmented_sort_by_key(table_view{{child}}, + table_view{{child}}, + output_offset->view(), + {column_order}, + {null_precedence}, + stream, + mr); + + return make_lists_column(input.size(), + std::move(output_offset), + std::move(sorted_child_table->release().front()), + input.null_count(), + cudf::detail::copy_bitmask(input.parent(), stream, mr)); +} } // namespace detail std::unique_ptr sort_lists(lists_column_view const& input, @@ -279,5 +313,15 @@ std::unique_ptr sort_lists(lists_column_view const& input, return detail::sort_lists(input, column_order, null_precedence, rmm::cuda_stream_default, mr); } +std::unique_ptr stable_sort_lists(lists_column_view const& input, + order column_order, + null_order null_precedence, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::stable_sort_lists( + input, column_order, null_precedence, rmm::cuda_stream_default, mr); +} + } // namespace lists } // namespace cudf diff --git a/cpp/src/sort/segmented_sort.cu b/cpp/src/sort/segmented_sort.cu index 9c4a2786612..8947da7e1bb 100644 --- a/cpp/src/sort/segmented_sort.cu +++ b/cpp/src/sort/segmented_sort.cu @@ -14,36 +14,26 @@ * limitations under the License. */ -#include -#include -#include -#include #include -#include #include #include -#include -#include -#include -#include -#include #include #include #include #include #include -#include - -#include -#include -#include -#include namespace cudf { namespace detail { +namespace { +/** + * @brief The enum specifying which sorting method to use (stable or unstable). + */ +enum class sort_method { STABLE, UNSTABLE }; + // returns segment indices for each element for all segments. // first segment begin index = 0, last segment end index = num_rows. rmm::device_uvector get_segment_indices(size_type num_rows, @@ -65,12 +55,14 @@ rmm::device_uvector get_segment_indices(size_type num_rows, return segment_ids; } -std::unique_ptr segmented_sorted_order(table_view const& keys, - column_view const& segment_offsets, - std::vector const& column_order, - std::vector const& null_precedence, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +std::unique_ptr segmented_sorted_order_common( + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order, + std::vector const& null_precedence, + sort_method sorting, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { CUDF_EXPECTS(segment_offsets.type() == data_type(type_to_id()), "segment offsets should be size_type"); @@ -95,26 +87,39 @@ std::unique_ptr segmented_sorted_order(table_view const& keys, }; auto child_column_order = prepend_default(column_order, order::ASCENDING); auto child_null_precedence = prepend_default(null_precedence, null_order::AFTER); + // return sorted order of child columns - return detail::sorted_order(segid_keys, child_column_order, child_null_precedence, stream, mr); + return sorting == sort_method::STABLE + ? detail::stable_sorted_order( + segid_keys, child_column_order, child_null_precedence, stream, mr) + : detail::sorted_order( + segid_keys, child_column_order, child_null_precedence, stream, mr); } -std::unique_ptr
segmented_sort_by_key(table_view const& values, - table_view const& keys, - column_view const& segment_offsets, - std::vector const& column_order, - std::vector const& null_precedence, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +std::unique_ptr
segmented_sort_by_key_common(table_view const& values, + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order, + std::vector const& null_precedence, + sort_method sorting, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { CUDF_EXPECTS(values.num_rows() == keys.num_rows(), "Mismatch in number of rows for values and keys"); - auto sorted_order = segmented_sorted_order(keys, - segment_offsets, - column_order, - null_precedence, - stream, - rmm::mr::get_current_device_resource()); + auto sorted_order = sorting == sort_method::STABLE + ? stable_segmented_sorted_order(keys, + segment_offsets, + column_order, + null_precedence, + stream, + rmm::mr::get_current_device_resource()) + : segmented_sorted_order(keys, + segment_offsets, + column_order, + null_precedence, + stream, + rmm::mr::get_current_device_resource()); // Gather segmented sort of child value columns` return detail::gather(values, @@ -124,8 +129,87 @@ std::unique_ptr
segmented_sort_by_key(table_view const& values, stream, mr); } + +} // namespace + +std::unique_ptr segmented_sorted_order(table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order, + std::vector const& null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + return segmented_sorted_order_common( + keys, segment_offsets, column_order, null_precedence, sort_method::UNSTABLE, stream, mr); +} + +std::unique_ptr stable_segmented_sorted_order( + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order, + std::vector const& null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + return segmented_sorted_order_common( + keys, segment_offsets, column_order, null_precedence, sort_method::STABLE, stream, mr); +} + +std::unique_ptr
segmented_sort_by_key(table_view const& values, + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order, + std::vector const& null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + return segmented_sort_by_key_common(values, + keys, + segment_offsets, + column_order, + null_precedence, + sort_method::UNSTABLE, + stream, + mr); +} + +std::unique_ptr
stable_segmented_sort_by_key(table_view const& values, + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order, + std::vector const& null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + return segmented_sort_by_key_common( + values, keys, segment_offsets, column_order, null_precedence, sort_method::STABLE, stream, mr); +} + } // namespace detail +std::unique_ptr segmented_sorted_order(table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order, + std::vector const& null_precedence, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::segmented_sorted_order( + keys, segment_offsets, column_order, null_precedence, rmm::cuda_stream_default, mr); +} + +std::unique_ptr stable_segmented_sorted_order( + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order, + std::vector const& null_precedence, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::stable_segmented_sorted_order( + keys, segment_offsets, column_order, null_precedence, rmm::cuda_stream_default, mr); +} + std::unique_ptr
segmented_sort_by_key(table_view const& values, table_view const& keys, column_view const& segment_offsets, @@ -138,4 +222,16 @@ std::unique_ptr
segmented_sort_by_key(table_view const& values, values, keys, segment_offsets, column_order, null_precedence, rmm::cuda_stream_default, mr); } +std::unique_ptr
stable_segmented_sort_by_key(table_view const& values, + table_view const& keys, + column_view const& segment_offsets, + std::vector const& column_order, + std::vector const& null_precedence, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + return detail::stable_segmented_sort_by_key( + values, keys, segment_offsets, column_order, null_precedence, rmm::cuda_stream_default, mr); +} + } // namespace cudf diff --git a/cpp/tests/lists/sort_lists_tests.cpp b/cpp/tests/lists/sort_lists_tests.cpp index d085d02d67b..ade626a5c2b 100644 --- a/cpp/tests/lists/sort_lists_tests.cpp +++ b/cpp/tests/lists/sort_lists_tests.cpp @@ -17,35 +17,29 @@ #include #include #include -#include #include -#include -#include -#include -#include #include -#include -#include -#include - -#include -#include -#include template using LCW = cudf::test::lists_column_wrapper; -using cudf::lists_column_view; -using cudf::lists::sort_lists; namespace cudf { namespace test { +auto generate_sorted_lists(lists_column_view const& input, + order column_order, + null_order null_precedence) +{ + return std::pair{lists::sort_lists(input, column_order, null_precedence), + lists::stable_sort_lists(input, column_order, null_precedence)}; +} + template struct SortLists : public BaseFixture { }; -TYPED_TEST_CASE(SortLists, NumericTypes); +TYPED_TEST_SUITE(SortLists, NumericTypes); using SortListsInt = SortLists; /* @@ -73,20 +67,34 @@ TYPED_TEST(SortLists, NoNull) // Ascending // LCW order{{2, 1, 0, 3}, {0}, {1, 2, 0}, {0, 1}}; LCW expected{{1, 2, 3, 4}, {5}, {8, 9, 10}, {6, 7}}; - auto results = sort_lists(lists_column_view{list}, order::ASCENDING, null_order::AFTER); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected); - - results = sort_lists(lists_column_view{list}, order::ASCENDING, null_order::BEFORE); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected); + { + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{list}, order::ASCENDING, null_order::AFTER); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), expected); + } + { + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{list}, order::ASCENDING, null_order::BEFORE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), expected); + } // Descending // LCW order{{3, 0, 1, 2}, {0}, {0, 1, 2}, {1, 0}}; LCW expected2{{4, 3, 2, 1}, {5}, {10, 9, 8}, {7, 6}}; - results = sort_lists(lists_column_view{list}, order::DESCENDING, null_order::AFTER); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected2); - - results = sort_lists(lists_column_view{list}, order::DESCENDING, null_order::BEFORE); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected2); + { + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{list}, order::DESCENDING, null_order::AFTER); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), expected2); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), expected2); + } + { + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{list}, order::DESCENDING, null_order::BEFORE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), expected2); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), expected2); + } } TYPED_TEST(SortLists, Null) @@ -100,60 +108,99 @@ TYPED_TEST(SortLists, Null) // List LCW list{{{3, 2, 4, 1}, valids_o.begin()}, {5}, {10, 8, 9}, {6, 7}}; // LCW order{{2, 1, 3, 0}, {0}, {1, 2, 0}, {0, 1}}; - LCW expected1{{{1, 2, 3, 4}, valids_a.begin()}, {5}, {8, 9, 10}, {6, 7}}; - LCW expected2{{{4, 1, 2, 3}, valids_b.begin()}, {5}, {8, 9, 10}, {6, 7}}; - auto results = sort_lists(lists_column_view{list}, order::ASCENDING, null_order::AFTER); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected1); - results = sort_lists(lists_column_view{list}, order::ASCENDING, null_order::BEFORE); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected2); + { + LCW expected{{{1, 2, 3, 4}, valids_a.begin()}, {5}, {8, 9, 10}, {6, 7}}; + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{list}, order::ASCENDING, null_order::AFTER); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), expected); + } + + { + LCW expected{{{4, 1, 2, 3}, valids_b.begin()}, {5}, {8, 9, 10}, {6, 7}}; + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{list}, order::ASCENDING, null_order::BEFORE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), expected); + } // Descending // LCW order{{3, 0, 1, 2}, {0}, {0, 1, 2}, {1, 0}}; - LCW expected3{{{4, 3, 2, 1}, valids_b.begin()}, {5}, {10, 9, 8}, {7, 6}}; - LCW expected4{{{3, 2, 1, 4}, valids_a.begin()}, {5}, {10, 9, 8}, {7, 6}}; - results = sort_lists(lists_column_view{list}, order::DESCENDING, null_order::AFTER); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected3); + { + LCW expected{{{4, 3, 2, 1}, valids_b.begin()}, {5}, {10, 9, 8}, {7, 6}}; + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{list}, order::DESCENDING, null_order::AFTER); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), expected); + } - results = sort_lists(lists_column_view{list}, order::DESCENDING, null_order::BEFORE); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), expected4); + { + LCW expected{{{3, 2, 1, 4}, valids_a.begin()}, {5}, {10, 9, 8}, {7, 6}}; + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{list}, order::DESCENDING, null_order::BEFORE); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), expected); + } } TEST_F(SortListsInt, Empty) { using T = int; - LCW l1{}; - LCW l2{LCW{}}; - LCW l3{LCW{}, LCW{}}; - - auto results = sort_lists(lists_column_view{l1}, {}, {}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), l1); - results = sort_lists(lists_column_view{l2}, {}, {}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), l2); - results = sort_lists(lists_column_view{l3}, {}, {}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), l3); + + { + LCW l{}; + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{l}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), l); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), l); + } + { + LCW l{LCW{}}; + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{l}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), l); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), l); + } + { + LCW l{LCW{}, LCW{}}; + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{l}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), l); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), l); + } } TEST_F(SortListsInt, Single) { using T = int; - LCW l1{{1}}; - LCW l2{{1, 2, 3}}; - auto results = sort_lists(lists_column_view{l1}, {}, {}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), l1); - results = sort_lists(lists_column_view{l2}, {}, {}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), l2); + { + LCW l{1}; + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{l}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), l); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), l); + } + { + LCW l{{1, 2, 3}}; + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{l}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), l); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), l); + } } TEST_F(SortListsInt, NullRows) { using T = int; std::vector valids{0, 1, 0}; - LCW l1{{{1, 2, 3}, {4, 5, 6}, {7}}, valids.begin()}; // offset 0, 0, 3, 3 + LCW l{{{1, 2, 3}, {4, 5, 6}, {7}}, valids.begin()}; // offset 0, 0, 3, 3 - auto results = sort_lists(lists_column_view{l1}, {}, {}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), l1); + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{l}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), l); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), l); } /* @@ -171,23 +218,43 @@ TEST_F(SortListsInt, Depth) TEST_F(SortListsInt, Sliced) { using T = int; - LCW l1{{3, 2, 1, 4}, {7, 5, 6}, {8, 9}, {10}}; + LCW l{{3, 2, 1, 4}, {7, 5, 6}, {8, 9}, {10}}; - auto sliced_list = cudf::slice(l1, {0, 4})[0]; - auto results = sort_lists(lists_column_view{sliced_list}, {}, {}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), LCW{{1, 2, 3, 4}, {5, 6, 7}, {8, 9}, {10}}); + { + auto const sliced_list = cudf::slice(l, {0, 4})[0]; + auto const expected = LCW{{1, 2, 3, 4}, {5, 6, 7}, {8, 9}, {10}}; + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{sliced_list}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), expected); + } - sliced_list = cudf::slice(l1, {1, 4})[0]; - results = sort_lists(lists_column_view{sliced_list}, {}, {}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), LCW{{5, 6, 7}, {8, 9}, {10}}); + { + auto const sliced_list = cudf::slice(l, {1, 4})[0]; + auto const expected = LCW{{5, 6, 7}, {8, 9}, {10}}; + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{sliced_list}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), expected); + } - sliced_list = cudf::slice(l1, {1, 2})[0]; - results = sort_lists(lists_column_view{sliced_list}, {}, {}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), LCW{{5, 6, 7}}); + { + auto const sliced_list = cudf::slice(l, {1, 2})[0]; + auto const expected = LCW{{5, 6, 7}}; + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{sliced_list}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), expected); + } - sliced_list = cudf::slice(l1, {0, 2})[0]; - results = sort_lists(lists_column_view{sliced_list}, {}, {}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(results->view(), LCW{{1, 2, 3, 4}, {5, 6, 7}}); + { + auto const sliced_list = cudf::slice(l, {0, 2})[0]; + auto const expected = LCW{{1, 2, 3, 4}, {5, 6, 7}}; + auto const [sorted_lists, stable_sorted_lists] = + generate_sorted_lists(lists_column_view{sliced_list}, {}, {}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorted_lists->view(), expected); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(stable_sorted_lists->view(), expected); + } } } // namespace test From d6acedd43e89338d8f32fea4338639525b62d175 Mon Sep 17 00:00:00 2001 From: MithunR Date: Mon, 18 Oct 2021 23:03:07 -0700 Subject: [PATCH 3/4] Push down parent nulls when flattening nested columns. (#9443) Fixes #9441. `cudf::structs::detail::flatten_nested_columns()` is used to flatten `STRUCT` columns, for use in operations where recursive descent is infeasible. E.g. JOINing tables on a `STRUCT` join key. If a `STRUCT` column's null-mask has not been superimposed over the member columns', it is possible that a null in a `STRUCT` row is not reflected in the corresponding member child/row. This causes the results of `flatten_nested_columns()` to cease to be equivalent to the `STRUCT` column. The incorrect output proceeds to corrupt the results of JOIN, as seen in #9441. The erstwhile convention (i.e. to call `superimpose_parent_nulls()` before `flatten_nested_columns()`) is easy to get wrong. This commit ensures that `superimpose_parent_nulls()` is called in `flatten_nested_columns()`. The previous explicit `superimpose` calls (e.g. in `rank_scan`) have been removed in favour of banking on `flatten_nested_columns()` doing this automatically. Authors: - MithunR (https://github.com/mythrocks) Approvers: - AJ Schmidt (https://github.com/ajschmidt8) - Jake Hemstad (https://github.com/jrhemstad) - Nghia Truong (https://github.com/ttnghia) URL: https://github.com/rapidsai/cudf/pull/9443 --- conda/recipes/libcudf/meta.yaml | 1 + .../cudf/detail/groupby/sort_helper.hpp | 3 +- .../cudf/detail}/structs/utilities.hpp | 109 ++++++++++++++++-- cpp/src/groupby/groupby.cu | 6 +- cpp/src/groupby/sort/group_rank_scan.cu | 7 +- cpp/src/groupby/sort/scan.cpp | 3 +- cpp/src/groupby/sort/sort_helper.cu | 10 +- cpp/src/join/hash_join.cu | 17 ++- cpp/src/join/hash_join.cuh | 2 + cpp/src/join/semi_join.cu | 6 +- cpp/src/lists/drop_list_duplicates.cu | 5 +- cpp/src/reductions/scan/rank_scan.cu | 7 +- cpp/src/scalar/scalar.cpp | 3 +- cpp/src/search/search.cu | 10 +- cpp/src/sort/is_sorted.cu | 8 +- cpp/src/sort/sort_impl.cuh | 14 +-- cpp/src/structs/copying/concatenate.cu | 2 +- cpp/src/structs/structs_column_factories.cu | 3 +- cpp/src/structs/utilities.cpp | 72 ++++++++---- cpp/tests/join/join_tests.cpp | 57 +++++++++ cpp/tests/structs/utilities_tests.cpp | 6 +- 21 files changed, 255 insertions(+), 96 deletions(-) rename cpp/{src => include/cudf/detail}/structs/utilities.hpp (58%) diff --git a/conda/recipes/libcudf/meta.yaml b/conda/recipes/libcudf/meta.yaml index e0b94efa210..8ccfdaa4aed 100644 --- a/conda/recipes/libcudf/meta.yaml +++ b/conda/recipes/libcudf/meta.yaml @@ -93,6 +93,7 @@ test: - test -f $PREFIX/include/cudf/detail/sequence.hpp - test -f $PREFIX/include/cudf/detail/sorting.hpp - test -f $PREFIX/include/cudf/detail/stream_compaction.hpp + - test -f $PREFIX/include/cudf/detail/structs/utilities.hpp - test -f $PREFIX/include/cudf/detail/tdigest/tdigest.hpp - test -f $PREFIX/include/cudf/detail/transform.hpp - test -f $PREFIX/include/cudf/detail/transpose.hpp diff --git a/cpp/include/cudf/detail/groupby/sort_helper.hpp b/cpp/include/cudf/detail/groupby/sort_helper.hpp index 1e36b2b2797..8705bbd29cb 100644 --- a/cpp/include/cudf/detail/groupby/sort_helper.hpp +++ b/cpp/include/cudf/detail/groupby/sort_helper.hpp @@ -18,6 +18,7 @@ #include #include +#include #include #include @@ -218,7 +219,7 @@ struct sort_groupby_helper { column_ptr _keys_bitmask_column; ///< Column representing rows with one or more nulls values table_view _keys; ///< Input keys to sort by table_view _unflattened_keys; ///< Input keys, unflattened and possibly nested - std::vector _struct_null_vectors; ///< Null vectors for struct columns in _keys + structs::detail::flattened_table _flattened; ///< Support datastructures for _keys index_vector_ptr _group_offsets; ///< Indices into sorted _keys indicating starting index of each groups diff --git a/cpp/src/structs/utilities.hpp b/cpp/include/cudf/detail/structs/utilities.hpp similarity index 58% rename from cpp/src/structs/utilities.hpp rename to cpp/include/cudf/detail/structs/utilities.hpp index 24b80b58669..aece79107c6 100644 --- a/cpp/src/structs/utilities.hpp +++ b/cpp/include/cudf/detail/structs/utilities.hpp @@ -16,10 +16,12 @@ #pragma once #include +#include #include #include #include +#include namespace cudf { namespace structs { @@ -64,6 +66,71 @@ std::vector> extract_ordered_struct_children( */ bool is_or_has_nested_lists(cudf::column_view const& col); +/** + * @brief Result of `flatten_nested_columns()`, where all `STRUCT` columns are replaced with + * their non-nested member columns, and `BOOL8` columns for their null masks. + * + * `flatten_nested_columns()` produces a "flattened" table_view with all `STRUCT` columns + * replaced with their child column_views, preceded by their null masks. + * All newly allocated columns and device_buffers that back the returned table_view + * are also encapsulated in `flatten_result`. + * + * Objects of `flatten_result` need to kept alive while its table_view is accessed. + */ +class flattened_table { + public: + /** + * @brief Constructor, to be used from `flatten_nested_columns()`. + * + * @param flattened_columns_ table_view resulting from `flatten_nested_columns()` + * @param orders_ Per-column ordering of the table_view + * @param null_orders_ Per-column null_order of the table_view + * @param columns_ Newly allocated columns to back the table_view + * @param null_masks_ Newly allocated null masks to back the table_view + */ + flattened_table(table_view const& flattened_columns_, + std::vector const& orders_, + std::vector const& null_orders_, + std::vector>&& columns_, + std::vector&& null_masks_) + : _flattened_columns{flattened_columns_}, + _orders{orders_}, + _null_orders{null_orders_}, + _columns{std::move(columns_)}, + _superimposed_nullmasks{std::move(null_masks_)} + { + } + + flattened_table() = default; + + /** + * @brief Getter for the flattened columns, as a `table_view`. + */ + table_view flattened_columns() const { return _flattened_columns; } + + /** + * @brief Getter for the cudf::order of the table_view's columns. + */ + std::vector orders() const { return _orders; } + + /** + * @brief Getter for the cudf::null_order of the table_view's columns. + */ + std::vector null_orders() const { return _null_orders; } + + /** + * @brief Conversion to `table_view`, to fetch flattened columns. + */ + operator table_view() const { return flattened_columns(); } + + private: + table_view _flattened_columns; + std::vector _orders; + std::vector _null_orders; + std::vector> _columns; + std::vector _superimposed_nullmasks; +}; + /** * @brief Flatten table with struct columns to table with constituent columns of struct columns. * @@ -74,17 +141,14 @@ bool is_or_has_nested_lists(cudf::column_view const& col); * @param null_precedence null order for input table * @param nullability force output to have nullability columns even if input columns * are all valid - * @return tuple with flattened table, flattened column order, flattened null precedence, - * vector of boolean columns (struct validity). + * @return `flatten_result` with flattened table, flattened column order, flattened null precedence, + * alongside the supporting columns and device_buffers for the flattened table. */ -std::tuple, - std::vector, - std::vector>> -flatten_nested_columns(table_view const& input, - std::vector const& column_order, - std::vector const& null_precedence, - column_nullability nullability = column_nullability::MATCH_INCOMING); +flattened_table flatten_nested_columns( + table_view const& input, + std::vector const& column_order, + std::vector const& null_precedence, + column_nullability nullability = column_nullability::MATCH_INCOMING); /** * @brief Unflatten columns flattened as by `flatten_nested_columns()`, @@ -156,6 +220,31 @@ std::tuple> superimpose_paren rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +/** + * @brief Push down nulls from a parent mask into child columns, using bitwise AND, + * for all columns in the specified table. + * + * This function constructs a table_view containing a new column_view instance equivalent to + * every column_view in the specified table. Each column_view might contain possibly new + * child column_views, all with possibly new null mask values reflecting null rows from the + * parent column: + * 1. If the column is not STRUCT, the column is returned unmodified, with no new + * supporting device_buffer instances. + * 2. If the column is STRUCT, the null masks of the parent and child are bitwise-ANDed, and a + * modified column_view is returned. This applies recursively. + * + * @param table The table_view of (possibly STRUCT) columns whose nulls need to be pushed to its + * members. + * @param stream CUDA stream used for device memory operations and kernel launches. + * @param mr Device memory resource used to allocate new device memory. + * @return A pair of: + * 1. table_view of columns with nulls pushed down to child columns, as appropriate. + * 2. Supporting device_buffer instances, for any newly constructed null masks. + */ +std::tuple> superimpose_parent_nulls( + table_view const& table, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); } // namespace detail } // namespace structs } // namespace cudf diff --git a/cpp/src/groupby/groupby.cu b/cpp/src/groupby/groupby.cu index bdaccba38dc..7ac7c199b05 100644 --- a/cpp/src/groupby/groupby.cu +++ b/cpp/src/groupby/groupby.cu @@ -25,6 +25,7 @@ #include #include #include +#include #include #include #include @@ -33,7 +34,6 @@ #include #include #include -#include #include @@ -76,8 +76,8 @@ std::pair, std::vector> groupby::disp if (_keys_are_sorted == sorted::NO and not _helper and detail::hash::can_use_hash_groupby(_keys, requests)) { // Optionally flatten nested key columns. - auto [flattened_keys, _, __, ___] = - flatten_nested_columns(_keys, {}, {}, column_nullability::FORCE); + auto flattened = flatten_nested_columns(_keys, {}, {}, column_nullability::FORCE); + auto flattened_keys = flattened.flattened_columns(); auto is_supported_key_type = [](auto col) { return cudf::is_equality_comparable(col.type()); }; CUDF_EXPECTS(std::all_of(flattened_keys.begin(), flattened_keys.end(), is_supported_key_type), "Unsupported groupby key type does not support equality comparison"); diff --git a/cpp/src/groupby/sort/group_rank_scan.cu b/cpp/src/groupby/sort/group_rank_scan.cu index 5f4dda294fd..df9f5b391fb 100644 --- a/cpp/src/groupby/sort/group_rank_scan.cu +++ b/cpp/src/groupby/sort/group_rank_scan.cu @@ -16,6 +16,7 @@ #include #include +#include #include #include #include @@ -23,8 +24,6 @@ #include #include -#include - namespace cudf { namespace groupby { namespace detail { @@ -55,9 +54,9 @@ std::unique_ptr rank_generator(column_view const& order_by, { auto const superimposed = structs::detail::superimpose_parent_nulls(order_by, stream, mr); table_view const order_table{{std::get<0>(superimposed)}}; - auto const flattener = cudf::structs::detail::flatten_nested_columns( + auto const flattened = cudf::structs::detail::flatten_nested_columns( order_table, {}, {}, structs::detail::column_nullability::MATCH_INCOMING); - auto const d_flat_order = table_device_view::create(std::get<0>(flattener), stream); + auto const d_flat_order = table_device_view::create(flattened, stream); row_equality_comparator comparator(*d_flat_order, *d_flat_order, true); auto ranks = make_fixed_width_column(data_type{type_to_id()}, order_table.num_rows(), diff --git a/cpp/src/groupby/sort/scan.cpp b/cpp/src/groupby/sort/scan.cpp index ace5d0e539c..b22f82ce7e4 100644 --- a/cpp/src/groupby/sort/scan.cpp +++ b/cpp/src/groupby/sort/scan.cpp @@ -22,6 +22,7 @@ #include #include #include +#include #include #include #include @@ -29,8 +30,6 @@ #include -#include - #include namespace cudf { diff --git a/cpp/src/groupby/sort/sort_helper.cu b/cpp/src/groupby/sort/sort_helper.cu index c4905b86ab9..1caf2ff0371 100644 --- a/cpp/src/groupby/sort/sort_helper.cu +++ b/cpp/src/groupby/sort/sort_helper.cu @@ -23,11 +23,11 @@ #include #include #include +#include #include #include #include #include -#include #include #include @@ -102,13 +102,11 @@ sort_groupby_helper::sort_groupby_helper(table_view const& keys, { using namespace cudf::structs::detail; - auto [flattened_keys, _, __, struct_null_vectors] = - flatten_nested_columns(keys, {}, {}, column_nullability::FORCE); + _flattened = flatten_nested_columns(keys, {}, {}, column_nullability::FORCE); + _keys = _flattened; auto is_supported_key_type = [](auto col) { return cudf::is_equality_comparable(col.type()); }; - CUDF_EXPECTS(std::all_of(flattened_keys.begin(), flattened_keys.end(), is_supported_key_type), + CUDF_EXPECTS(std::all_of(_keys.begin(), _keys.end(), is_supported_key_type), "Unsupported groupby key type does not support equality comparison"); - _struct_null_vectors = std::move(struct_null_vectors); - _keys = flattened_keys; // Cannot depend on caller's sorting if the column contains nulls, // and null values are to be excluded. diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu index 07ad2e052f1..c1a5f8f17c3 100644 --- a/cpp/src/join/hash_join.cu +++ b/cpp/src/join/hash_join.cu @@ -15,11 +15,11 @@ */ #include #include -#include #include #include #include +#include #include #include @@ -293,11 +293,10 @@ hash_join::hash_join_impl::hash_join_impl(cudf::table_view const& build, CUDF_EXPECTS(build.num_rows() < cudf::detail::MAX_JOIN_SIZE, "Build column size is too big for hash join"); - auto flattened_build = structs::detail::flatten_nested_columns( - build, {}, {}, structs::detail::column_nullability::FORCE); - _build = std::get<0>(flattened_build); // need to store off the owning structures for some of the views in _build - _created_null_columns = std::move(std::get<3>(flattened_build)); + _flattened_build_table = structs::detail::flatten_nested_columns( + build, {}, {}, structs::detail::column_nullability::FORCE); + _build = _flattened_build_table; if (0 == build.num_rows()) { return; } @@ -354,7 +353,7 @@ std::size_t hash_join::hash_join_impl::inner_join_size(cudf::table_view const& p auto flattened_probe = structs::detail::flatten_nested_columns( probe, {}, {}, structs::detail::column_nullability::FORCE); - auto const flattened_probe_table = std::get<0>(flattened_probe); + auto const flattened_probe_table = flattened_probe.flattened_columns(); auto build_table_ptr = cudf::table_device_view::create(_build, stream); auto flattened_probe_table_ptr = cudf::table_device_view::create(flattened_probe_table, stream); @@ -374,7 +373,7 @@ std::size_t hash_join::hash_join_impl::left_join_size(cudf::table_view const& pr auto flattened_probe = structs::detail::flatten_nested_columns( probe, {}, {}, structs::detail::column_nullability::FORCE); - auto const flattened_probe_table = std::get<0>(flattened_probe); + auto const flattened_probe_table = flattened_probe.flattened_columns(); auto build_table_ptr = cudf::table_device_view::create(_build, stream); auto flattened_probe_table_ptr = cudf::table_device_view::create(flattened_probe_table, stream); @@ -395,7 +394,7 @@ std::size_t hash_join::hash_join_impl::full_join_size(cudf::table_view const& pr auto flattened_probe = structs::detail::flatten_nested_columns( probe, {}, {}, structs::detail::column_nullability::FORCE); - auto const flattened_probe_table = std::get<0>(flattened_probe); + auto const flattened_probe_table = flattened_probe.flattened_columns(); auto build_table_ptr = cudf::table_device_view::create(_build, stream); auto flattened_probe_table_ptr = cudf::table_device_view::create(flattened_probe_table, stream); @@ -419,7 +418,7 @@ hash_join::hash_join_impl::compute_hash_join(cudf::table_view const& probe, auto flattened_probe = structs::detail::flatten_nested_columns( probe, {}, {}, structs::detail::column_nullability::FORCE); - auto const flattened_probe_table = std::get<0>(flattened_probe); + auto const flattened_probe_table = flattened_probe.flattened_columns(); CUDF_EXPECTS(_build.num_columns() == flattened_probe_table.num_columns(), "Mismatch in number of columns to be joined on"); diff --git a/cpp/src/join/hash_join.cuh b/cpp/src/join/hash_join.cuh index 8951d8b3aca..1d39daed457 100644 --- a/cpp/src/join/hash_join.cuh +++ b/cpp/src/join/hash_join.cuh @@ -20,6 +20,7 @@ #include #include +#include #include #include #include @@ -138,6 +139,7 @@ struct hash_join::hash_join_impl { private: cudf::table_view _build; std::vector> _created_null_columns; + cudf::structs::detail::flattened_table _flattened_build_table; std::unique_ptr> _hash_table; diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu index 69a7b8c722b..4a2f46d6f43 100644 --- a/cpp/src/join/semi_join.cu +++ b/cpp/src/join/semi_join.cu @@ -16,12 +16,12 @@ #include #include -#include #include #include #include #include +#include #include #include #include @@ -68,8 +68,8 @@ std::unique_ptr> left_semi_anti_join( auto left_flattened_tables = structs::detail::flatten_nested_columns( left_keys, {}, {}, structs::detail::column_nullability::FORCE); - auto right_flattened_keys = std::get<0>(right_flattened_tables); - auto left_flattened_keys = std::get<0>(left_flattened_tables); + auto right_flattened_keys = right_flattened_tables.flattened_columns(); + auto left_flattened_keys = left_flattened_tables.flattened_columns(); // Only care about existence, so we'll use an unordered map (other joins need a multimap) using hash_table_type = concurrent_unordered_map; diff --git a/cpp/src/lists/drop_list_duplicates.cu b/cpp/src/lists/drop_list_duplicates.cu index c547ca14f2d..c9ebcca183d 100644 --- a/cpp/src/lists/drop_list_duplicates.cu +++ b/cpp/src/lists/drop_list_duplicates.cu @@ -14,14 +14,13 @@ * limitations under the License. */ -#include - #include #include #include #include #include #include +#include #include #include #include @@ -487,7 +486,7 @@ struct get_unique_entries_dispatch { : structs::detail::column_nullability::MATCH_INCOMING; auto const entries_flattened = cudf::structs::detail::flatten_nested_columns( entries_tview, {order::ASCENDING}, {null_order::AFTER}, flatten_nullability); - auto const d_view = table_device_view::create(std::get<0>(entries_flattened), stream); + auto const d_view = table_device_view::create(entries_flattened, stream); auto const comp = table_row_comparator_fn{list_offsets, *d_view, diff --git a/cpp/src/reductions/scan/rank_scan.cu b/cpp/src/reductions/scan/rank_scan.cu index 566b9aadea8..bb6a85094f5 100644 --- a/cpp/src/reductions/scan/rank_scan.cu +++ b/cpp/src/reductions/scan/rank_scan.cu @@ -14,10 +14,9 @@ * limitations under the License. */ -#include - #include #include +#include #include #include @@ -53,9 +52,9 @@ std::unique_ptr rank_generator(column_view const& order_by, { auto const superimposed = structs::detail::superimpose_parent_nulls(order_by, stream, mr); table_view const order_table{{std::get<0>(superimposed)}}; - auto const flattener = cudf::structs::detail::flatten_nested_columns( + auto const flattened = cudf::structs::detail::flatten_nested_columns( order_table, {}, {}, structs::detail::column_nullability::MATCH_INCOMING); - auto const d_flat_order = table_device_view::create(std::get<0>(flattener), stream); + auto const d_flat_order = table_device_view::create(flattened, stream); row_equality_comparator comparator(*d_flat_order, *d_flat_order, true); auto ranks = make_fixed_width_column(data_type{type_to_id()}, order_table.num_rows(), diff --git a/cpp/src/scalar/scalar.cpp b/cpp/src/scalar/scalar.cpp index f982e7b99f2..5b7abdfcaf0 100644 --- a/cpp/src/scalar/scalar.cpp +++ b/cpp/src/scalar/scalar.cpp @@ -14,10 +14,9 @@ * limitations under the License. */ -#include - #include #include +#include #include #include #include diff --git a/cpp/src/search/search.cu b/cpp/src/search/search.cu index 705fd1a2372..462d0678eab 100644 --- a/cpp/src/search/search.cu +++ b/cpp/src/search/search.cu @@ -18,6 +18,7 @@ #include #include #include +#include #include #include #include @@ -26,7 +27,6 @@ #include #include #include -#include #include @@ -112,13 +112,13 @@ std::unique_ptr search_ordered(table_view const& t, auto const values_flattened = structs::detail::flatten_nested_columns(matched.second.back(), {}, {}, flatten_nullability); - auto const t_d = table_device_view::create(std::get<0>(t_flattened), stream); - auto const values_d = table_device_view::create(std::get<0>(values_flattened), stream); + auto const t_d = table_device_view::create(t_flattened, stream); + auto const values_d = table_device_view::create(values_flattened, stream); auto const& lhs = find_first ? *t_d : *values_d; auto const& rhs = find_first ? *values_d : *t_d; - auto const& column_order_flattened = std::get<1>(t_flattened); - auto const& null_precedence_flattened = std::get<2>(t_flattened); + auto const& column_order_flattened = t_flattened.orders(); + auto const& null_precedence_flattened = t_flattened.null_orders(); auto const column_order_dv = detail::make_device_uvector_async(column_order_flattened, stream); auto const null_precedence_dv = detail::make_device_uvector_async(null_precedence_flattened, stream); diff --git a/cpp/src/sort/is_sorted.cu b/cpp/src/sort/is_sorted.cu index c9ab791395d..b08baaa0261 100644 --- a/cpp/src/sort/is_sorted.cu +++ b/cpp/src/sort/is_sorted.cu @@ -15,13 +15,13 @@ */ #include +#include #include #include #include #include #include #include -#include #include #include @@ -39,10 +39,10 @@ auto is_sorted(cudf::table_view const& in, // 0-table_view, 1-column_order, 2-null_precedence, 3-validity_columns auto flattened = structs::detail::flatten_nested_columns(in, column_order, null_precedence); - auto const d_input = table_device_view::create(std::get<0>(flattened), stream); - auto const d_column_order = make_device_uvector_async(std::get<1>(flattened), stream); + auto const d_input = table_device_view::create(flattened, stream); + auto const d_column_order = make_device_uvector_async(flattened.orders(), stream); auto const d_null_precedence = has_nulls - ? make_device_uvector_async(std::get<2>(flattened), stream) + ? make_device_uvector_async(flattened.null_orders(), stream) : rmm::device_uvector(0, stream); auto comparator = row_lexicographic_comparator( diff --git a/cpp/src/sort/sort_impl.cuh b/cpp/src/sort/sort_impl.cuh index d707ece5ba9..25f0815e645 100644 --- a/cpp/src/sort/sort_impl.cuh +++ b/cpp/src/sort/sort_impl.cuh @@ -18,14 +18,13 @@ #include #include +#include #include #include #include #include #include -#include - #include #include #include @@ -124,13 +123,12 @@ 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& input_flattened = std::get<0>(flattened); - auto device_table = table_device_view::create(input_flattened, stream); - auto const d_column_order = make_device_uvector_async(std::get<1>(flattened), stream); + 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); - if (has_nulls(input_flattened)) { - auto const d_null_precedence = make_device_uvector_async(std::get<2>(flattened), stream); + if (has_nulls(flattened)) { + auto const d_null_precedence = make_device_uvector_async(flattened.null_orders(), stream); auto const comparator = row_lexicographic_comparator( *device_table, *device_table, d_column_order.data(), d_null_precedence.data()); if (stable) { diff --git a/cpp/src/structs/copying/concatenate.cu b/cpp/src/structs/copying/concatenate.cu index fe5483b119d..7656470b791 100644 --- a/cpp/src/structs/copying/concatenate.cu +++ b/cpp/src/structs/copying/concatenate.cu @@ -21,8 +21,8 @@ #include #include #include +#include #include -#include #include diff --git a/cpp/src/structs/structs_column_factories.cu b/cpp/src/structs/structs_column_factories.cu index 833ceab7518..ee401c82bcb 100644 --- a/cpp/src/structs/structs_column_factories.cu +++ b/cpp/src/structs/structs_column_factories.cu @@ -14,9 +14,8 @@ * limitations under the License. */ -#include - #include +#include #include #include diff --git a/cpp/src/structs/utilities.cpp b/cpp/src/structs/utilities.cpp index b84af73b681..47f8f29385c 100644 --- a/cpp/src/structs/utilities.cpp +++ b/cpp/src/structs/utilities.cpp @@ -14,10 +14,9 @@ * limitations under the License. */ -#include - #include #include +#include #include #include #include @@ -26,9 +25,13 @@ #include #include #include -#include + +#include + +#include #include +#include namespace cudf { namespace structs { @@ -85,30 +88,40 @@ bool is_or_has_nested_lists(cudf::column_view const& col) * @brief Flattens struct columns to constituent non-struct columns in the input table. * */ -struct flattened_table { +struct table_flattener { + table_view input; // reference variables - table_view const& input; std::vector const& column_order; std::vector const& null_precedence; // output std::vector> validity_as_column; + std::vector superimposed_nullmasks; std::vector flat_columns; std::vector flat_column_order; std::vector flat_null_precedence; column_nullability nullability; - flattened_table(table_view const& input, + table_flattener(table_view const& input, std::vector const& column_order, std::vector const& null_precedence, column_nullability nullability) - : input(input), - column_order(column_order), - null_precedence(null_precedence), - nullability(nullability) + : column_order(column_order), null_precedence(null_precedence), nullability(nullability) { + superimpose_nulls(input); fail_if_unsupported_types(input); } + /** + * @brief Pushes down nulls from struct columns to children, saves the resulting + * column to `input`, and generated null masks to `superimposed_nullmasks`. + */ + void superimpose_nulls(table_view const& input_table) + { + auto [table, null_masks] = superimpose_parent_nulls(input_table); + this->input = table; + this->superimposed_nullmasks = std::move(null_masks); + } + void fail_if_unsupported_types(table_view const& input) const { auto const has_lists = std::any_of(input.begin(), input.end(), is_or_has_nested_lists); @@ -176,29 +189,23 @@ struct flattened_table { } } - return std::make_tuple(table_view{flat_columns}, + return flattened_table{table_view{flat_columns}, std::move(flat_column_order), std::move(flat_null_precedence), - std::move(validity_as_column)); + std::move(validity_as_column), + std::move(superimposed_nullmasks)}; } }; -std::tuple, - std::vector, - std::vector>> -flatten_nested_columns(table_view const& input, - std::vector const& column_order, - std::vector const& null_precedence, - column_nullability nullability) +flattened_table flatten_nested_columns(table_view const& input, + std::vector const& column_order, + std::vector const& null_precedence, + column_nullability nullability) { auto const has_struct = std::any_of(input.begin(), input.end(), is_struct); - if (not has_struct) { - return std::make_tuple( - input, column_order, null_precedence, std::vector>{}); - } + if (not has_struct) { return flattened_table{input, column_order, null_precedence, {}, {}}; } - return flattened_table{input, column_order, null_precedence, nullability}(); + return table_flattener{input, column_order, null_precedence, nullability}(); } namespace { @@ -415,6 +422,21 @@ std::tuple> superimpose_paren std::move(ret_validity_buffers)); } +std::tuple> superimpose_parent_nulls( + table_view const& table, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) +{ + auto superimposed_columns = std::vector{}; + auto superimposed_nullmasks = std::vector{}; + for (auto col : table) { + auto [superimposed_col, null_masks] = superimpose_parent_nulls(col); + superimposed_columns.push_back(superimposed_col); + superimposed_nullmasks.insert(superimposed_nullmasks.begin(), + std::make_move_iterator(null_masks.begin()), + std::make_move_iterator(null_masks.end())); + } + return {table_view{superimposed_columns}, std::move(superimposed_nullmasks)}; +} + } // namespace detail } // namespace structs } // namespace cudf diff --git a/cpp/tests/join/join_tests.cpp b/cpp/tests/join/join_tests.cpp index 8945f82baef..6a515efc3bb 100644 --- a/cpp/tests/join/join_tests.cpp +++ b/cpp/tests/join/join_tests.cpp @@ -18,6 +18,8 @@ #include #include #include +#include +#include #include #include #include @@ -30,6 +32,7 @@ #include #include #include +#include #include #include @@ -1735,4 +1738,58 @@ TEST_F(JoinTest, FullJoinWithStructsAndNulls) CUDF_TEST_EXPECT_TABLES_EQUIVALENT(*sorted_gold, *sorted_result); } +TEST_F(JoinTest, Repro_StructsWithoutNullsPushedDown) +{ + // When joining on a STRUCT column, if the parent nulls are not reflected in + // the children, the join might produce incorrect results. + // + // In this test, a fact table of structs is joined against a dimension table. + // Both tables must match (only) on the NULL row. This will fail if the fact table's + // nulls are not pushed down into its children. + using ints = column_wrapper; + using structs = cudf::test::structs_column_wrapper; + using namespace cudf::test::iterators; + + auto make_table = [](auto&& col) { + auto columns = CVector{}; + columns.push_back(std::move(col)); + return cudf::table{std::move(columns)}; + }; + + auto const fact_table = [make_table] { + auto fact_ints = ints{0, 1, 2, 3, 4}; + auto fact_structs = structs{{fact_ints}, no_nulls()}.release(); + // Now set struct validity to invalidate index#3. + cudf::detail::set_null_mask(fact_structs->mutable_view().null_mask(), 3, 4, false); + // Struct row#3 is null, but Struct.child has a non-null value. + return make_table(std::move(fact_structs)); + }(); + + auto const dimension_table = [make_table] { + auto dim_ints = ints{999}; + auto dim_structs = structs{{dim_ints}, null_at(0)}; + return make_table(dim_structs.release()); + }(); + + auto const result = cudf::inner_join(fact_table.view(), dimension_table.view(), {0}, {0}); + EXPECT_EQ(result->num_rows(), 1); // The null STRUCT rows should match. + + // Note: Join result might not have nulls pushed down, since it's an output of gather(). + // Must superimpose parent nulls before comparisons. + auto [superimposed_results, _] = cudf::structs::detail::superimpose_parent_nulls(*result); + + auto const expected = [] { + auto fact_ints = ints{0}; + auto fact_structs = structs{{fact_ints}, null_at(0)}; + auto dim_ints = ints{0}; + auto dim_structs = structs{{dim_ints}, null_at(0)}; + auto columns = CVector{}; + columns.push_back(fact_structs.release()); + columns.push_back(dim_structs.release()); + return cudf::table{std::move(columns)}; + }(); + + CUDF_TEST_EXPECT_TABLES_EQUIVALENT(superimposed_results, expected); +} + CUDF_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/structs/utilities_tests.cpp b/cpp/tests/structs/utilities_tests.cpp index 08b40b22aa4..b26ea87c5b8 100644 --- a/cpp/tests/structs/utilities_tests.cpp +++ b/cpp/tests/structs/utilities_tests.cpp @@ -14,8 +14,6 @@ * limitations under the License. */ -#include - #include #include #include @@ -27,6 +25,7 @@ #include #include #include +#include #include namespace cudf::test { @@ -39,8 +38,7 @@ void flatten_unflatten_compare(table_view const& input_table) { using namespace cudf::structs::detail; - auto [flattened, _, __, ___] = - flatten_nested_columns(input_table, {}, {}, column_nullability::FORCE); + auto flattened = flatten_nested_columns(input_table, {}, {}, column_nullability::FORCE); auto unflattened = unflatten_nested_columns(std::make_unique(flattened), input_table); From a3a27c69faa9712b5ab7b8bd5b5e2561efee9666 Mon Sep 17 00:00:00 2001 From: Karthikeyan <6488848+karthikeyann@users.noreply.github.com> Date: Tue, 19 Oct 2021 13:33:07 +0530 Subject: [PATCH 4/4] add ctest memcheck using cuda-sanitizer (#9414) addresses part of https://github.com/rapidsai/cudf/issues/904 and pre-requisite for https://github.com/rapidsai/ops/issues/1462 - Adds `cuda-sanitizer` as memcheck tool in ctest. - Adds environment variable `GTEST_CUDF_RMM_MODE` in gtests fixture This PR enables the unit tests to run under `compute-sanitizer --tool memcheck` in gpuCI nightly or weekly builds. The main issue pending in this PR is to find suitable way for specifying **"--rmm_mode=cuda"** only in memcheck runs. `ctest` doesn't allow passing arguments to tests yet ([ctest feature request](https://gitlab.kitware.com/cmake/cmake/-/issues/20470)). So, options are 1. environment variable in cmake command (requires cmake rebuild_cache!) - `export GTEST_CUDF_RMM_MODE=cuda; ninja rebuild_cache; ctest -T memcheck` 2. environment variable in ctest using [RunTests script](https://stackoverflow.com/a/54675911/1550940) (has limitations and fails sometimes) 3. ctest -S custom_script 4. add custom target in cmake instead of inbuilt ctest memcheck 5. Add environment variable in gtest for rmm mode. (Another easier option) *(This is implemented in this PR)* - `export GTEST_CUDF_RMM_MODE=cuda; ctest -T memcheck` Which method is most suitable? implemented [5] Authors: - Karthikeyan (https://github.com/karthikeyann) Approvers: - https://github.com/nvdbaranec - Robert Maynard (https://github.com/robertmaynard) - David Wendt (https://github.com/davidwendt) URL: https://github.com/rapidsai/cudf/pull/9414 --- cpp/CMakeLists.txt | 5 +++++ cpp/include/cudf_test/base_fixture.hpp | 10 ++++++++-- cpp/tests/CMakeLists.txt | 2 +- 3 files changed, 14 insertions(+), 3 deletions(-) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 3bfdf23c2bc..8aed7089dc5 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -97,6 +97,11 @@ rapids_find_package(CUDAToolkit REQUIRED INSTALL_EXPORT_SET cudf-exports) include(cmake/Modules/ConfigureCUDA.cmake) # set other CUDA compilation flags +# ctest cuda memcheck +find_program(CUDA_SANITIZER compute-sanitizer) +set(MEMORYCHECK_COMMAND ${CUDA_SANITIZER}) +set(MEMORYCHECK_TYPE CudaSanitizer) +set(CUDA_SANITIZER_COMMAND_OPTIONS "--tool memcheck") ################################################################################################### # - dependencies ---------------------------------------------------------------------------------- diff --git a/cpp/include/cudf_test/base_fixture.hpp b/cpp/include/cudf_test/base_fixture.hpp index 6f69dc19cfd..eee2d396d9f 100644 --- a/cpp/include/cudf_test/base_fixture.hpp +++ b/cpp/include/cudf_test/base_fixture.hpp @@ -266,6 +266,9 @@ inline std::shared_ptr create_memory_resource( * * Currently only supports 'rmm_mode' string parameter, which set the rmm * allocation mode. The default value of the parameter is 'pool'. + * Environment variable 'CUDF_TEST_RMM_MODE' can also be used to set the rmm + * allocation mode. If both are set, the value of 'rmm_mode' string parameter + * takes precedence. * * @return Parsing results in the form of unordered map */ @@ -273,9 +276,12 @@ inline auto parse_cudf_test_opts(int argc, char** argv) { try { cxxopts::Options options(argv[0], " - cuDF tests command line options"); + const char* env_rmm_mode = std::getenv("GTEST_CUDF_RMM_MODE"); // Overridden by CLI options + auto default_rmm_mode = env_rmm_mode ? env_rmm_mode : "pool"; options.allow_unrecognised_options().add_options()( - "rmm_mode", "RMM allocation mode", cxxopts::value()->default_value("pool")); - + "rmm_mode", + "RMM allocation mode", + cxxopts::value()->default_value(default_rmm_mode)); return options.parse(argc, argv); } catch (const cxxopts::OptionException& e) { CUDF_FAIL("Error parsing command line options"); diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index f61bafe5133..eb4e8735638 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -387,7 +387,7 @@ ConfigureTest(STRINGS_TEST ################################################################################################### # - structs test ---------------------------------------------------------------------------------- -ConfigureTest(STRUCTS_TEST +ConfigureTest(STRUCTS_TEST structs/structs_column_tests.cpp structs/utilities_tests.cpp )