From 4e7336bcce56b418b0c6e387f1005e7b6d73b26f Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Tue, 13 Apr 2021 16:47:27 -0600 Subject: [PATCH 01/36] Add `lists/interleave_columns.cuh|cu` files --- cpp/CMakeLists.txt | 3 ++- cpp/include/cudf/lists/detail/interleave_columns.cuh | 8 ++++++++ cpp/src/lists/interleave_columns.cu | 3 +++ 3 files changed, 13 insertions(+), 1 deletion(-) create mode 100644 cpp/include/cudf/lists/detail/interleave_columns.cuh create mode 100644 cpp/src/lists/interleave_columns.cu diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 98f6ed40502..a299c872de8 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -265,9 +265,10 @@ add_library(cudf src/lists/copying/gather.cu src/lists/copying/segmented_gather.cu src/lists/count_elements.cu + src/lists/drop_list_duplicates.cu src/lists/explode.cu src/lists/extract.cu - src/lists/drop_list_duplicates.cu + src/lists/interleave_columns.cu src/lists/lists_column_factories.cu src/lists/lists_column_view.cu src/lists/segmented_sort.cu diff --git a/cpp/include/cudf/lists/detail/interleave_columns.cuh b/cpp/include/cudf/lists/detail/interleave_columns.cuh new file mode 100644 index 00000000000..7e07b92e3d1 --- /dev/null +++ b/cpp/include/cudf/lists/detail/interleave_columns.cuh @@ -0,0 +1,8 @@ +// +// Created by Nghia Truong on 4/13/21. +// + +#ifndef CUDF_INTERLEAVE_COLUMNS_CUH +#define CUDF_INTERLEAVE_COLUMNS_CUH + +#endif // CUDF_INTERLEAVE_COLUMNS_CUH diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu new file mode 100644 index 00000000000..346d14dfc2c --- /dev/null +++ b/cpp/src/lists/interleave_columns.cu @@ -0,0 +1,3 @@ +// +// Created by Nghia Truong on 4/13/21. +// From 4921a8d808906c164ef7c0e871861d81695f62fc Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Tue, 13 Apr 2021 19:33:55 -0600 Subject: [PATCH 02/36] Remove headers from `interleave_columns_tests.cpp` --- cpp/tests/reshape/interleave_columns_tests.cpp | 6 ------ 1 file changed, 6 deletions(-) diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index 654df7589e6..d507fad1cf7 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -16,17 +16,11 @@ #include #include -#include #include -#include #include -#include #include -#include -#include "cudf/utilities/traits.hpp" - using namespace cudf::test; template From e61170ec52afe69a7924666b0f70c62185362c02 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Tue, 13 Apr 2021 19:49:39 -0600 Subject: [PATCH 03/36] Add skeleton for `interleave_columns` that supports `list_view` --- cpp/src/reshape/interleave_columns.cu | 16 ++++++++++++++-- 1 file changed, 14 insertions(+), 2 deletions(-) diff --git a/cpp/src/reshape/interleave_columns.cu b/cpp/src/reshape/interleave_columns.cu index 3e2cb4ac02f..f7c0ae4be11 100644 --- a/cpp/src/reshape/interleave_columns.cu +++ b/cpp/src/reshape/interleave_columns.cu @@ -30,11 +30,23 @@ namespace detail { namespace { struct interleave_columns_functor { template - std::enable_if_t() and not std::is_same::value, + std::enable_if_t() and + not std::is_same::value and + not std::is_same::value, std::unique_ptr> operator()(Args&&... args) { - CUDF_FAIL("interleave_columns not supported for dictionary and list types."); + CUDF_FAIL("Called `interleave_columns` on none-supported data type."); + } + + template + std::enable_if_t::value, std::unique_ptr> + operator()(table_view const& strings_columns, + bool create_mask, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + return std::make_unique(*(strings_columns.begin()), stream, mr); } template From 07beb16c4beb502670f1baccd0200134b316a82b Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 14 Apr 2021 07:39:03 -0600 Subject: [PATCH 04/36] Finish doxygen and declaration for `lists::interleave_columns` --- .../cudf/lists/detail/interleave_columns.cuh | 8 --- .../cudf/lists/detail/interleave_columns.hpp | 53 +++++++++++++++++++ 2 files changed, 53 insertions(+), 8 deletions(-) delete mode 100644 cpp/include/cudf/lists/detail/interleave_columns.cuh create mode 100644 cpp/include/cudf/lists/detail/interleave_columns.hpp diff --git a/cpp/include/cudf/lists/detail/interleave_columns.cuh b/cpp/include/cudf/lists/detail/interleave_columns.cuh deleted file mode 100644 index 7e07b92e3d1..00000000000 --- a/cpp/include/cudf/lists/detail/interleave_columns.cuh +++ /dev/null @@ -1,8 +0,0 @@ -// -// Created by Nghia Truong on 4/13/21. -// - -#ifndef CUDF_INTERLEAVE_COLUMNS_CUH -#define CUDF_INTERLEAVE_COLUMNS_CUH - -#endif // CUDF_INTERLEAVE_COLUMNS_CUH diff --git a/cpp/include/cudf/lists/detail/interleave_columns.hpp b/cpp/include/cudf/lists/detail/interleave_columns.hpp new file mode 100644 index 00000000000..1bb4993ae0b --- /dev/null +++ b/cpp/include/cudf/lists/detail/interleave_columns.hpp @@ -0,0 +1,53 @@ +/* + * Copyright (c) 2021, 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 + +namespace cudf { +namespace lists { +namespace detail { + +/** + * @brief Returns a single column by interleaving rows of the given table of list elements. + * + * @code{.pseudo} + * s1 = [{0, 1}, {2, 3, 4}, {5}, {}, {6, 7}] + * s2 = [{8}, {9}, {}, {10, 11, 12}, {13, 14, 15, 16}] + * r = interleave_columns(s1, s2) + * r is now [{0, 1}, {8}, {2, 3, 4}, {9}, {5}, {}, {}, {10, 11, 12}, {6, 7}, {13, 14, 15, 16}] + * @endcode + * + * @throws cudf::logic_error if columns of the input table are not lists columns. + * + * @param input Table containing columns to interleave. + * @param stream CUDA stream used for device memory operations and kernel launches. + * @param mr Device memory resource used to allocate the returned column's device + * memory. + * @return The interleaved columns as a single column. + */ +std::unique_ptr interleave_columns( + table_view const& lists_columns, + bool create_mask, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + +} // namespace detail +} // namespace lists +} // namespace cudf From 90f69736408e5ba61627cebccc5a4007b4820d9b Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 14 Apr 2021 07:53:10 -0600 Subject: [PATCH 05/36] Add call to `lists::interleave_columns` --- cpp/src/reshape/interleave_columns.cu | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/cpp/src/reshape/interleave_columns.cu b/cpp/src/reshape/interleave_columns.cu index f7c0ae4be11..f5c1c3f3147 100644 --- a/cpp/src/reshape/interleave_columns.cu +++ b/cpp/src/reshape/interleave_columns.cu @@ -19,6 +19,7 @@ #include #include #include +#include #include #include @@ -41,12 +42,12 @@ struct interleave_columns_functor { template std::enable_if_t::value, std::unique_ptr> - operator()(table_view const& strings_columns, + operator()(table_view const& lists_columns, bool create_mask, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - return std::make_unique(*(strings_columns.begin()), stream, mr); + return lists::detail::interleave_columns(lists_columns, create_mask, stream, mr); } template From f70fb411aac29bab040fedfbdfb57c7fe881c4e4 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 14 Apr 2021 09:19:25 -0600 Subject: [PATCH 06/36] WIP for `lists::interleave_columns` --- cpp/src/lists/interleave_columns.cu | 98 ++++++++++++++++++++++++++++- 1 file changed, 95 insertions(+), 3 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 346d14dfc2c..7e82d81252c 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -1,3 +1,95 @@ -// -// Created by Nghia Truong on 4/13/21. -// +/* + * Copyright (c) 2021, 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 + +#include +#include +#include + +#include + +namespace cudf { +namespace lists { +namespace detail { +namespace { + +} // anonymous namespace + +/** + * @copydoc cudf::lists::detail::interleave_columns + * + */ +std::unique_ptr interleave_columns(table_view const& lists_columns, + bool create_mask, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_EXPECTS(std::all_of(lists_columns.begin(), + lists_columns.end(), + [](auto const& col) { return col.type().id() == type_id::LIST; }), + "All columns must be of type list"); + CUDF_EXPECTS(std::all_of(lists_columns.begin(), + lists_columns.end(), + [](auto const& col) { + return not cudf::is_nested(lists_column_view(col).child().type()); + }), + "Nested types are not supported in `lists::interleave_columns`"); + + // Single column returns a copy + if (lists_columns.num_columns() == 1) { + return std::make_unique(*(lists_columns.begin()), stream, mr); + } + + auto const num_rows = lists_columns.num_rows(); + if (num_rows == 0) { return cudf::empty_like(lists_columns.column(0)); } + + static_assert(sizeof(offset_type) == sizeof(int32_t)); + auto lists_offsets = make_numeric_column( + data_type{type_id::INT32}, num_rows + 1, mask_state::UNALLOCATED, stream, mr); + auto const output_offsets_ptr = lists_offsets->mutable_view().begin(); + + auto lists_entries = std::make_unique(); + + static constexpr auto invalid_size = std::numeric_limits::lowest(); + auto const count_it = thrust::make_counting_iterator(0); + auto [null_mask, null_count] = cudf::detail::valid_if( + count_it, + count_it + num_rows, + [str_sizes = output_offsets_ptr + 1] __device__(size_type idx) { + return str_sizes[idx] != invalid_size; + }, + stream, + mr); + + return make_lists_column(num_rows, + std::move(lists_offsets), + std::move(lists_entries), + null_count, + null_count ? std::move(null_mask) : rmm::device_buffer{}, + stream, + mr); +} + +} // namespace detail +} // namespace lists +} // namespace cudf From 7aa2120b5139caf0002ac0d7bd8a249e7bfc9181 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 14 Apr 2021 09:42:38 -0600 Subject: [PATCH 07/36] Finish InvalidInput test --- .../reshape/interleave_columns_tests.cpp | 55 +++++++++++++++++++ 1 file changed, 55 insertions(+) diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index d507fad1cf7..61d43c77c96 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -368,4 +368,59 @@ TYPED_TEST(FixedPointTestBothReps, FixedPointInterleave) } } +namespace { +using StrListsCol = cudf::test::lists_column_wrapper; +using IntListsCol = cudf::test::lists_column_wrapper; +using IntCol = cudf::test::fixed_width_column_wrapper; +using TView = cudf::table_view; + +constexpr bool print_all{true}; // For debugging +constexpr int32_t null{0}; + +auto null_at(cudf::size_type idx) +{ + return cudf::detail::make_counting_transform_iterator(0, [idx](auto i) { return i != idx; }); +} +auto null_at(std::vector const& indices) +{ + return cudf::detail::make_counting_transform_iterator(0, [&indices](auto i) { + return std::find(indices.cbegin(), indices.cend(), i) == indices.cend(); + }); +} + +auto all_nulls() +{ + return cudf::detail::make_counting_transform_iterator(0, [](auto) { return false; }); +} + +} // namespace + +struct ListsColumnsInterleaveTest : public cudf::test::BaseFixture { +}; + +TEST_F(ListsColumnsInterleaveTest, InvalidInput) +{ + // Input table contains non-list column + { + auto const col1 = IntCol{}.release(); + auto const col2 = IntListsCol{}.release(); + EXPECT_THROW(cudf::interleave_columns(TView{{col1->view(), col2->view()}}), cudf::logic_error); + } + + // Nested types are not supported + { + auto const col = IntListsCol{{IntListsCol{1, 2, 3}, IntListsCol{4, 5, 6}}}.release(); + EXPECT_THROW(cudf::interleave_columns(TView{{col->view(), col->view()}}), cudf::logic_error); + } +} + +template +struct ListsColumnsInterleaveTypedTest : public cudf::test::BaseFixture { +}; +#define ListsCol cudf::test::lists_column_wrapper + +using TypesForTest = + cudf::test::Concat; +TYPED_TEST_CASE(ListsColumnsInterleaveTypedTest, TypesForTest); + CUDF_TEST_PROGRAM_MAIN() From 3065577189d0a229f4d865c8c0b81e5dc4d19a35 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 14 Apr 2021 10:03:09 -0600 Subject: [PATCH 08/36] Rewrite input validation --- cpp/src/lists/interleave_columns.cu | 21 +++++++++++---------- 1 file changed, 11 insertions(+), 10 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 7e82d81252c..bc0466aa713 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -44,16 +44,17 @@ std::unique_ptr interleave_columns(table_view const& lists_columns, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - CUDF_EXPECTS(std::all_of(lists_columns.begin(), - lists_columns.end(), - [](auto const& col) { return col.type().id() == type_id::LIST; }), - "All columns must be of type list"); - CUDF_EXPECTS(std::all_of(lists_columns.begin(), - lists_columns.end(), - [](auto const& col) { - return not cudf::is_nested(lists_column_view(col).child().type()); - }), - "Nested types are not supported in `lists::interleave_columns`"); + auto const first_col_entry_type_id = + lists_column_view(*lists_columns.begin()).child().type().id(); + for (auto const& col : lists_columns) { + CUDF_EXPECTS(col.type().id() == type_id::LIST, + "All columns of the input table must be of lists column type."); + + auto const child_col = lists_column_view(col).child(); + CUDF_EXPECTS(not cudf::is_nested(child_col.type()), "Nested types are not supported."); + CUDF_EXPECTS(first_col_entry_type_id == child_col.type().id(), + "The types of entries in the input columns must be the same."); + } // Single column returns a copy if (lists_columns.num_columns() == 1) { From 180cc215c50d7f2dcd16680a3d492338a2e846db Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 14 Apr 2021 10:03:30 -0600 Subject: [PATCH 09/36] Rewrite InvalidInput test and add skeleton for other tests --- .../reshape/interleave_columns_tests.cpp | 266 ++++++++++++++++++ 1 file changed, 266 insertions(+) diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index 61d43c77c96..13d08694dee 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -407,6 +407,13 @@ TEST_F(ListsColumnsInterleaveTest, InvalidInput) EXPECT_THROW(cudf::interleave_columns(TView{{col1->view(), col2->view()}}), cudf::logic_error); } + // Types mismatch + { + auto const col1 = IntListsCol{}.release(); + auto const col2 = StrListsCol{}.release(); + EXPECT_THROW(cudf::interleave_columns(TView{{col1->view(), col2->view()}}), cudf::logic_error); + } + // Nested types are not supported { auto const col = IntListsCol{{IntListsCol{1, 2, 3}, IntListsCol{4, 5, 6}}}.release(); @@ -423,4 +430,263 @@ using TypesForTest = cudf::test::Concat; TYPED_TEST_CASE(ListsColumnsInterleaveTypedTest, TypesForTest); +TYPED_TEST(ListsColumnsInterleaveTypedTest, ConcatenateEmptyColumns) +{ + auto const col1 = ListsCol{}.release(); + auto const col2 = ListsCol{}.release(); + auto const expected = ListsCol{}.release(); + auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view()}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); +} + +TYPED_TEST(ListsColumnsInterleaveTypedTest, ConcatenateOneColumnNotNull) +{ + auto const col = ListsCol{{1, 2}, {3, 4}, {5, 6}}.release(); + auto const results = cudf::interleave_columns(TView{{col->view()}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*col, *results, print_all); +} + +TYPED_TEST(ListsColumnsInterleaveTypedTest, ConcatenateOneColumnWithNulls) +{ + auto const col = ListsCol{{ListsCol{{1, 2, null}, null_at(2)}, + ListsCol{} /*NULL*/, + ListsCol{{null, 3, 4, 4, 4, 4}, null_at(0)}, + ListsCol{5, 6}}, + null_at(1)} + .release(); + auto const results = cudf::interleave_columns(TView{{col->view()}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*col, *results, print_all); +} + +TYPED_TEST(ListsColumnsInterleaveTypedTest, SimpleInputNoNull) +{ + auto const col1 = ListsCol{{1, 2}, {3, 4}, {5, 6}}.release(); + auto const col2 = ListsCol{{7, 8}, {9, 10}, {11, 12}}.release(); + auto const expected = ListsCol{{1, 2, 7, 8}, {3, 4, 9, 10}, {5, 6, 11, 12}}.release(); + auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view()}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); +} + +TEST_F(ListsColumnsInterleaveTest, SimpleInputStringsColumnsNoNull) +{ + auto const col1 = StrListsCol{ + StrListsCol{"Tomato", "Apple"}, + StrListsCol{"Banana", "Kiwi", "Cherry"}, + StrListsCol{ + "Coconut"}}.release(); + auto const col2 = + StrListsCol{StrListsCol{"Orange"}, StrListsCol{"Lemon", "Peach"}, StrListsCol{}}.release(); + auto const expected = StrListsCol{ + StrListsCol{"Tomato", "Apple", "Orange"}, + StrListsCol{"Banana", "Kiwi", "Cherry", "Lemon", "Peach"}, + StrListsCol{ + "Coconut"}}.release(); + auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view()}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); +} + +TYPED_TEST(ListsColumnsInterleaveTypedTest, SimpleInputWithNulls) +{ + auto const col1 = ListsCol{{ListsCol{{1, null, 3, 4}, null_at(1)}, + ListsCol{{null, 2, 3, 4}, null_at(0)}, + ListsCol{{null, 2, 3, 4}, null_at(0)}, + ListsCol{} /*NULL*/, + ListsCol{{1, 2, null, 4}, null_at(2)}, + ListsCol{{1, 2, 3, null}, null_at(3)}}, + null_at(3)} + .release(); + auto const col2 = ListsCol{{ListsCol{{10, 11, 12, null}, null_at(3)}, + ListsCol{{13, 14, 15, 16, 17, null}, null_at(5)}, + ListsCol{} /*NULL*/, + ListsCol{{null, 18}, null_at(0)}, + ListsCol{{19, 20, null}, null_at(2)}, + ListsCol{{null}, null_at(0)}}, + null_at(2)} + .release(); + auto const col3 = ListsCol{{ListsCol{} /*NULL*/, + ListsCol{{20, null}, null_at(1)}, + ListsCol{{null, 21, null, null}, null_at({0, 2, 3})}, + ListsCol{}, + ListsCol{22, 23, 24, 25}, + ListsCol{{null, null, null, null, null}, all_nulls()}}, + null_at(0)} + .release(); + + // Ignore null list elements + { + auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view()}}); + auto const expected = + ListsCol{{ListsCol{{1, null, 3, 4, 10, 11, 12, null}, null_at({1, 7})}, + ListsCol{{null, 2, 3, 4, 13, 14, 15, 16, 17, null, 20, null}, null_at({0, 9, 11})}, + ListsCol{{null, 2, 3, 4, null, 21, null, null}, null_at({0, 4, 6, 7})}, + ListsCol{{null, 18}, null_at(0)}, + ListsCol{{1, 2, null, 4, 19, 20, null, 22, 23, 24, 25}, null_at({2, 6})}, + ListsCol{{1, 2, 3, null, null, null, null, null, null, null}, + null_at({3, 4, 5, 6, 7, 8, 9})}}} + .release(); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); + } + + // Null list rows result in null list rows + { + auto const results = + cudf::interleave_columns(TView{{col1->view(), col2->view()}}, + cudf::lists::concatenate_null_policy::NULLIFY_OUTPUT_ROW); + auto const expected = + ListsCol{{ListsCol{} /*NULL*/, + ListsCol{{null, 2, 3, 4, 13, 14, 15, 16, 17, null, 20, null}, null_at({0, 9, 11})}, + ListsCol{} /*NULL*/, + ListsCol{} /*NULL*/, + ListsCol{{1, 2, null, 4, 19, 20, null, 22, 23, 24, 25}, null_at({2, 6})}, + ListsCol{{1, 2, 3, null, null, null, null, null, null, null}, + null_at({3, 4, 5, 6, 7, 8, 9})}}, + null_at({0, 2, 3})} + .release(); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); + } +} + +TEST_F(ListsColumnsInterleaveTest, SimpleInputStringsColumnsWithNulls) +{ + auto const col1 = StrListsCol{ + StrListsCol{{"Tomato", "Bear" /*NULL*/, "Apple"}, null_at(1)}, + StrListsCol{{"Banana", "Pig" /*NULL*/, "Kiwi", "Cherry", "Whale" /*NULL*/}, null_at({1, 4})}, + StrListsCol{ + "Coconut"}}.release(); + auto const col2 = + StrListsCol{ + {StrListsCol{{"Orange", "Dog" /*NULL*/, "Fox" /*NULL*/, "Duck" /*NULL*/}, null_at({1, 2, 3})}, + StrListsCol{"Lemon", "Peach"}, + StrListsCol{{"Deer" /*NULL*/, "Snake" /*NULL*/, "Horse" /*NULL*/}, all_nulls()}}, /*NULL*/ + null_at(2)} + .release(); + + // Ignore null list elements + { + auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view()}}); + auto const expected = StrListsCol{ + StrListsCol{{"Tomato", + "Bear" /*NULL*/, + "Apple", + "Orange", + "Dog" /*NULL*/, + "Fox" /*NULL*/, + "Duck" /*NULL*/}, + null_at({1, 4, 5, 6})}, + StrListsCol{{"Banana", "Pig" /*NULL*/, "Kiwi", "Cherry", "Whale" /*NULL*/, "Lemon", "Peach"}, + null_at({1, 4})}, + StrListsCol{ + "Coconut"}}.release(); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); + } + + // Null list rows result in null list rows + { + auto const results = + cudf::interleave_columns(TView{{col1->view(), col2->view()}}, + cudf::lists::concatenate_null_policy::NULLIFY_OUTPUT_ROW); + auto const expected = + StrListsCol{ + {StrListsCol{{"Tomato", + "Bear" /*NULL*/, + "Apple", + "Orange", + "Dog" /*NULL*/, + "Fox" /*NULL*/, + "Duck" /*NULL*/}, + null_at({1, 4, 5, 6})}, + StrListsCol{ + {"Banana", "Pig" /*NULL*/, "Kiwi", "Cherry", "Whale" /*NULL*/, "Lemon", "Peach"}, + null_at({1, 4})}, + StrListsCol{""} /*NULL*/}, + null_at(2)} + .release(); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); + } +} + +TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputNoNull) +{ + auto const col_original = ListsCol{{1, 2, 3}, {2, 3}, {3, 4, 5, 6}, {5, 6}, {}, {7}}.release(); + + // Slice first half + { + auto const col = cudf::slice(col_original->view(), {0, 3})[0]; + auto const expected = + ListsCol{ + {1, 2, 3, 1, 2, 3, 1, 2, 3}, {2, 3, 2, 3, 2, 3}, {3, 4, 5, 6, 3, 4, 5, 6, 3, 4, 5, 6}} + .release(); + auto const results = cudf::interleave_columns(TView{{col, col, col}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); + } + + // Slice second half + { + auto const col = cudf::slice(col_original->view(), {3, 6})[0]; + auto const expected = ListsCol{{5, 6, 5, 6, 5, 6}, {}, {7, 7, 7}}.release(); + auto const results = cudf::interleave_columns(TView{{col, col, col}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); + } + + // Slice middle segment + { + auto const col = cudf::slice(col_original->view(), {2, 4})[0]; + auto const expected = + ListsCol{{3, 4, 5, 6, 3, 4, 5, 6, 3, 4, 5, 6}, {5, 6, 5, 6, 5, 6}}.release(); + auto const results = cudf::interleave_columns(TView{{col, col, col}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); + } +} + +TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputWithNulls) +{ + auto const col_original = ListsCol{{ListsCol{{null, 2, 3}, null_at(0)}, + ListsCol{{2, 3}}, /*NULL*/ + ListsCol{{3, null, 5, 6}, null_at(1)}, + ListsCol{{5, 6}}, /*NULL*/ + ListsCol{}, /*NULL*/ + ListsCol{7}, + ListsCol{8, 9, 10}}, + null_at({1, 3, 4})} + .release(); + + // Slice first half + { + auto const col = cudf::slice(col_original->view(), {0, 3})[0]; + auto const expected = ListsCol{{ListsCol{{null, 2, 3, null, 2, 3}, null_at({0, 3})}, + ListsCol{}, /*NULL*/ + ListsCol{{3, null, 5, 6, 3, null, 5, 6}, null_at({1, 5})}}, + null_at(1)} + .release(); + auto const results = cudf::interleave_columns(TView{{col, col}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); + } + + // Slice second half + { + auto const col = cudf::slice(col_original->view(), {3, 7})[0]; + auto const expected = ListsCol{{ListsCol{}, /*NULL*/ + ListsCol{}, /*NULL*/ + ListsCol{7, 7}, + ListsCol{8, 9, 10, 8, 9, 10}}, + null_at({0, 1})} + .release(); + auto const results = cudf::interleave_columns(TView{{col, col}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); + } + + // Slice middle segment + { + auto const col = cudf::slice(col_original->view(), {2, 6})[0]; + auto const expected = ListsCol{{ListsCol{{3, null, 5, 6, 3, null, 5, 6}, null_at({1, 5})}, + ListsCol{}, /*NULL*/ + ListsCol{}, /*NULL*/ + ListsCol{7, 7}}, + null_at({1, 2})} + .release(); + auto const results = cudf::interleave_columns(TView{{col, col}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); + } +} + CUDF_TEST_PROGRAM_MAIN() From b7f994e5618d7a90d339f56972a6cc1a82b8501c Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 14 Apr 2021 11:48:07 -0600 Subject: [PATCH 10/36] Rewrite ConcatenateEmptyColumns test --- cpp/tests/reshape/interleave_columns_tests.cpp | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index 13d08694dee..1a67c38a8b1 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -432,11 +432,9 @@ TYPED_TEST_CASE(ListsColumnsInterleaveTypedTest, TypesForTest); TYPED_TEST(ListsColumnsInterleaveTypedTest, ConcatenateEmptyColumns) { - auto const col1 = ListsCol{}.release(); - auto const col2 = ListsCol{}.release(); - auto const expected = ListsCol{}.release(); - auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view()}}); - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); + auto const col = ListsCol{}.release(); + auto const results = cudf::interleave_columns(TView{{col->view(), col->view()}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*col, *results, print_all); } TYPED_TEST(ListsColumnsInterleaveTypedTest, ConcatenateOneColumnNotNull) From c5a58149d93439dab94f389b26bae15b5bae81a6 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 14 Apr 2021 12:10:11 -0600 Subject: [PATCH 11/36] Finish SimpleInputNoNull and SimpleInputStringsColumnsNoNull tests --- cpp/tests/reshape/interleave_columns_tests.cpp | 18 ++++++++++-------- 1 file changed, 10 insertions(+), 8 deletions(-) diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index 1a67c38a8b1..5df91298775 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -430,21 +430,21 @@ using TypesForTest = cudf::test::Concat; TYPED_TEST_CASE(ListsColumnsInterleaveTypedTest, TypesForTest); -TYPED_TEST(ListsColumnsInterleaveTypedTest, ConcatenateEmptyColumns) +TYPED_TEST(ListsColumnsInterleaveTypedTest, InterleaveEmptyColumns) { auto const col = ListsCol{}.release(); auto const results = cudf::interleave_columns(TView{{col->view(), col->view()}}); CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*col, *results, print_all); } -TYPED_TEST(ListsColumnsInterleaveTypedTest, ConcatenateOneColumnNotNull) +TYPED_TEST(ListsColumnsInterleaveTypedTest, InterleaveOneColumnNotNull) { auto const col = ListsCol{{1, 2}, {3, 4}, {5, 6}}.release(); auto const results = cudf::interleave_columns(TView{{col->view()}}); CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*col, *results, print_all); } -TYPED_TEST(ListsColumnsInterleaveTypedTest, ConcatenateOneColumnWithNulls) +TYPED_TEST(ListsColumnsInterleaveTypedTest, InterleaveOneColumnWithNulls) { auto const col = ListsCol{{ListsCol{{1, 2, null}, null_at(2)}, ListsCol{} /*NULL*/, @@ -460,7 +460,7 @@ TYPED_TEST(ListsColumnsInterleaveTypedTest, SimpleInputNoNull) { auto const col1 = ListsCol{{1, 2}, {3, 4}, {5, 6}}.release(); auto const col2 = ListsCol{{7, 8}, {9, 10}, {11, 12}}.release(); - auto const expected = ListsCol{{1, 2, 7, 8}, {3, 4, 9, 10}, {5, 6, 11, 12}}.release(); + auto const expected = ListsCol{{1, 2}, {7, 8}, {3, 4}, {9, 10}, {5, 6}, {11, 12}}.release(); auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view()}}); CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); } @@ -475,10 +475,12 @@ TEST_F(ListsColumnsInterleaveTest, SimpleInputStringsColumnsNoNull) auto const col2 = StrListsCol{StrListsCol{"Orange"}, StrListsCol{"Lemon", "Peach"}, StrListsCol{}}.release(); auto const expected = StrListsCol{ - StrListsCol{"Tomato", "Apple", "Orange"}, - StrListsCol{"Banana", "Kiwi", "Cherry", "Lemon", "Peach"}, - StrListsCol{ - "Coconut"}}.release(); + StrListsCol{"Tomato", "Apple"}, + StrListsCol{"Orange"}, + StrListsCol{"Banana", "Kiwi", "Cherry"}, + StrListsCol{"Lemon", "Peach"}, + StrListsCol{"Coconut"}, + StrListsCol{}}.release(); auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view()}}); CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); } From ea08f4391a5ec757c8ce10dee0f7e39f7a9f2cda Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 14 Apr 2021 12:28:30 -0600 Subject: [PATCH 12/36] Finish SimpleInputWithNulls test --- .../reshape/interleave_columns_tests.cpp | 54 ++++++++----------- 1 file changed, 22 insertions(+), 32 deletions(-) diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index 5df91298775..941743f65f8 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -512,38 +512,28 @@ TYPED_TEST(ListsColumnsInterleaveTypedTest, SimpleInputWithNulls) null_at(0)} .release(); - // Ignore null list elements - { - auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view()}}); - auto const expected = - ListsCol{{ListsCol{{1, null, 3, 4, 10, 11, 12, null}, null_at({1, 7})}, - ListsCol{{null, 2, 3, 4, 13, 14, 15, 16, 17, null, 20, null}, null_at({0, 9, 11})}, - ListsCol{{null, 2, 3, 4, null, 21, null, null}, null_at({0, 4, 6, 7})}, - ListsCol{{null, 18}, null_at(0)}, - ListsCol{{1, 2, null, 4, 19, 20, null, 22, 23, 24, 25}, null_at({2, 6})}, - ListsCol{{1, 2, 3, null, null, null, null, null, null, null}, - null_at({3, 4, 5, 6, 7, 8, 9})}}} - .release(); - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); - } - - // Null list rows result in null list rows - { - auto const results = - cudf::interleave_columns(TView{{col1->view(), col2->view()}}, - cudf::lists::concatenate_null_policy::NULLIFY_OUTPUT_ROW); - auto const expected = - ListsCol{{ListsCol{} /*NULL*/, - ListsCol{{null, 2, 3, 4, 13, 14, 15, 16, 17, null, 20, null}, null_at({0, 9, 11})}, - ListsCol{} /*NULL*/, - ListsCol{} /*NULL*/, - ListsCol{{1, 2, null, 4, 19, 20, null, 22, 23, 24, 25}, null_at({2, 6})}, - ListsCol{{1, 2, 3, null, null, null, null, null, null, null}, - null_at({3, 4, 5, 6, 7, 8, 9})}}, - null_at({0, 2, 3})} - .release(); - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); - } + auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view(), col3->view()}}); + auto const expected = ListsCol{{ListsCol{{1, null, 3, 4}, null_at(1)}, + ListsCol{{10, 11, 12, null}, null_at(3)}, + ListsCol{} /*NULL*/, + ListsCol{{null, 2, 3, 4}, null_at(0)}, + ListsCol{{13, 14, 15, 16, 17, null}, null_at(5)}, + ListsCol{{20, null}, null_at(1)}, + ListsCol{{null, 2, 3, 4}, null_at(0)}, + ListsCol{} /*NULL*/, + ListsCol{{null, 21, null, null}, null_at({0, 2, 3})}, + ListsCol{} /*NULL*/, + ListsCol{{null, 18}, null_at(0)}, + ListsCol{}, + ListsCol{{1, 2, null, 4}, null_at(2)}, + ListsCol{{19, 20, null}, null_at(2)}, + ListsCol{22, 23, 24, 25}, + ListsCol{{1, 2, 3, null}, null_at(3)}, + ListsCol{{null}, null_at(0)}, + ListsCol{{null, null, null, null, null}, all_nulls()}}, + null_at({2, 7, 9})} + .release(); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); } TEST_F(ListsColumnsInterleaveTest, SimpleInputStringsColumnsWithNulls) From 83305be5750eb05e95ee144898fd8b46e97b8e6f Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 14 Apr 2021 13:04:47 -0600 Subject: [PATCH 13/36] Finish SimpleInputStringsColumnsWithNulls --- .../reshape/interleave_columns_tests.cpp | 54 +++++-------------- 1 file changed, 12 insertions(+), 42 deletions(-) diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index 941743f65f8..7c75d51d971 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -551,48 +551,18 @@ TEST_F(ListsColumnsInterleaveTest, SimpleInputStringsColumnsWithNulls) null_at(2)} .release(); - // Ignore null list elements - { - auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view()}}); - auto const expected = StrListsCol{ - StrListsCol{{"Tomato", - "Bear" /*NULL*/, - "Apple", - "Orange", - "Dog" /*NULL*/, - "Fox" /*NULL*/, - "Duck" /*NULL*/}, - null_at({1, 4, 5, 6})}, - StrListsCol{{"Banana", "Pig" /*NULL*/, "Kiwi", "Cherry", "Whale" /*NULL*/, "Lemon", "Peach"}, - null_at({1, 4})}, - StrListsCol{ - "Coconut"}}.release(); - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); - } - - // Null list rows result in null list rows - { - auto const results = - cudf::interleave_columns(TView{{col1->view(), col2->view()}}, - cudf::lists::concatenate_null_policy::NULLIFY_OUTPUT_ROW); - auto const expected = - StrListsCol{ - {StrListsCol{{"Tomato", - "Bear" /*NULL*/, - "Apple", - "Orange", - "Dog" /*NULL*/, - "Fox" /*NULL*/, - "Duck" /*NULL*/}, - null_at({1, 4, 5, 6})}, - StrListsCol{ - {"Banana", "Pig" /*NULL*/, "Kiwi", "Cherry", "Whale" /*NULL*/, "Lemon", "Peach"}, - null_at({1, 4})}, - StrListsCol{""} /*NULL*/}, - null_at(2)} - .release(); - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); - } + auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view()}}); + auto const expected = + StrListsCol{ + {StrListsCol{{"Tomato", "" /*NULL*/, "Apple"}, null_at(1)}, + StrListsCol{{"Orange", "" /*NULL*/, "" /*NULL*/, "" /*NULL*/}, null_at({1, 2, 3})}, + StrListsCol{{"Banana", "" /*NULL*/, "Kiwi", "Cherry", "" /*NULL*/}, null_at({1, 4})}, + StrListsCol{"Lemon", "Peach"}, + StrListsCol{"Coconut"}, + StrListsCol{}}, /*NULL*/ + null_at(5)} + .release(); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); } TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputNoNull) From 76c487634d57d063e253f9a8ae97208ce000c90c Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 14 Apr 2021 14:04:55 -0600 Subject: [PATCH 14/36] Finish SlicedColumnsInputNoNull and SlicedColumnsInputWithNulls tests --- .../reshape/interleave_columns_tests.cpp | 112 +++++++----------- 1 file changed, 45 insertions(+), 67 deletions(-) diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index 7c75d51d971..56f29033842 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -512,7 +512,6 @@ TYPED_TEST(ListsColumnsInterleaveTypedTest, SimpleInputWithNulls) null_at(0)} .release(); - auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view(), col3->view()}}); auto const expected = ListsCol{{ListsCol{{1, null, 3, 4}, null_at(1)}, ListsCol{{10, 11, 12, null}, null_at(3)}, ListsCol{} /*NULL*/, @@ -533,6 +532,7 @@ TYPED_TEST(ListsColumnsInterleaveTypedTest, SimpleInputWithNulls) ListsCol{{null, null, null, null, null}, all_nulls()}}, null_at({2, 7, 9})} .release(); + auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view(), col3->view()}}); CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); } @@ -551,7 +551,6 @@ TEST_F(ListsColumnsInterleaveTest, SimpleInputStringsColumnsWithNulls) null_at(2)} .release(); - auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view()}}); auto const expected = StrListsCol{ {StrListsCol{{"Tomato", "" /*NULL*/, "Apple"}, null_at(1)}, @@ -562,40 +561,32 @@ TEST_F(ListsColumnsInterleaveTest, SimpleInputStringsColumnsWithNulls) StrListsCol{}}, /*NULL*/ null_at(5)} .release(); + auto const results = cudf::interleave_columns(TView{{col1->view(), col2->view()}}); CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); } TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputNoNull) { auto const col_original = ListsCol{{1, 2, 3}, {2, 3}, {3, 4, 5, 6}, {5, 6}, {}, {7}}.release(); - - // Slice first half - { - auto const col = cudf::slice(col_original->view(), {0, 3})[0]; - auto const expected = - ListsCol{ - {1, 2, 3, 1, 2, 3, 1, 2, 3}, {2, 3, 2, 3, 2, 3}, {3, 4, 5, 6, 3, 4, 5, 6, 3, 4, 5, 6}} - .release(); - auto const results = cudf::interleave_columns(TView{{col, col, col}}); - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); - } - - // Slice second half - { - auto const col = cudf::slice(col_original->view(), {3, 6})[0]; - auto const expected = ListsCol{{5, 6, 5, 6, 5, 6}, {}, {7, 7, 7}}.release(); - auto const results = cudf::interleave_columns(TView{{col, col, col}}); - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); - } - - // Slice middle segment - { - auto const col = cudf::slice(col_original->view(), {2, 4})[0]; - auto const expected = - ListsCol{{3, 4, 5, 6, 3, 4, 5, 6, 3, 4, 5, 6}, {5, 6, 5, 6, 5, 6}}.release(); - auto const results = cudf::interleave_columns(TView{{col, col, col}}); - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); - } + auto const col1 = cudf::slice(col_original->view(), {0, 3})[0]; + auto const col2 = cudf::slice(col_original->view(), {1, 4})[0]; + auto const col3 = cudf::slice(col_original->view(), {2, 5})[0]; + auto const col4 = cudf::slice(col_original->view(), {3, 6})[0]; + auto const expected = ListsCol{ + ListsCol{1, 2, 3}, + ListsCol{2, 3}, + ListsCol{3, 4, 5, 6}, + ListsCol{5, 6}, + ListsCol{2, 3}, + ListsCol{3, 4, 5, 6}, + ListsCol{5, 6}, + ListsCol{}, + ListsCol{3, 4, 5, 6}, + ListsCol{5, 6}, + ListsCol{}, + ListsCol{7}}.release(); + auto const results = cudf::interleave_columns(TView{{col1, col2, col3, col4}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); } TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputWithNulls) @@ -610,43 +601,30 @@ TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputWithNulls) null_at({1, 3, 4})} .release(); - // Slice first half - { - auto const col = cudf::slice(col_original->view(), {0, 3})[0]; - auto const expected = ListsCol{{ListsCol{{null, 2, 3, null, 2, 3}, null_at({0, 3})}, - ListsCol{}, /*NULL*/ - ListsCol{{3, null, 5, 6, 3, null, 5, 6}, null_at({1, 5})}}, - null_at(1)} - .release(); - auto const results = cudf::interleave_columns(TView{{col, col}}); - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); - } - - // Slice second half - { - auto const col = cudf::slice(col_original->view(), {3, 7})[0]; - auto const expected = ListsCol{{ListsCol{}, /*NULL*/ - ListsCol{}, /*NULL*/ - ListsCol{7, 7}, - ListsCol{8, 9, 10, 8, 9, 10}}, - null_at({0, 1})} - .release(); - auto const results = cudf::interleave_columns(TView{{col, col}}); - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); - } - - // Slice middle segment - { - auto const col = cudf::slice(col_original->view(), {2, 6})[0]; - auto const expected = ListsCol{{ListsCol{{3, null, 5, 6, 3, null, 5, 6}, null_at({1, 5})}, - ListsCol{}, /*NULL*/ - ListsCol{}, /*NULL*/ - ListsCol{7, 7}}, - null_at({1, 2})} - .release(); - auto const results = cudf::interleave_columns(TView{{col, col}}); - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); - } + auto const col1 = cudf::slice(col_original->view(), {0, 3})[0]; + auto const col2 = cudf::slice(col_original->view(), {1, 4})[0]; + auto const col3 = cudf::slice(col_original->view(), {2, 5})[0]; + auto const col4 = cudf::slice(col_original->view(), {3, 6})[0]; + auto const col5 = cudf::slice(col_original->view(), {4, 7})[0]; + auto const expected = ListsCol{{ListsCol{{null, 2, 3}, null_at(0)}, + ListsCol{}, /*NULL*/ + ListsCol{{3, null, 5, 6}, null_at(1)}, + ListsCol{}, /*NULL*/ + ListsCol{}, /*NULL*/ + ListsCol{}, /*NULL*/ + ListsCol{{3, null, 5, 6}, null_at(1)}, + ListsCol{}, /*NULL*/ + ListsCol{}, /*NULL*/ + ListsCol{7}, + ListsCol{{3, null, 5, 6}, null_at(1)}, + ListsCol{}, /*NULL*/ + ListsCol{}, /*NULL*/ + ListsCol{7}, + ListsCol{8, 9, 10}}, + null_at({1, 3, 4, 5, 7, 8, 11, 12})} + .release(); + auto const results = cudf::interleave_columns(TView{{col1, col2, col3, col4, col5}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); } CUDF_TEST_PROGRAM_MAIN() From 414dc9b076bae4f80961085bbbeb4dd3aa3b8f0c Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 21 Apr 2021 10:56:50 -0600 Subject: [PATCH 15/36] Rewrite doxygen --- cpp/include/cudf/lists/detail/interleave_columns.hpp | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/cpp/include/cudf/lists/detail/interleave_columns.hpp b/cpp/include/cudf/lists/detail/interleave_columns.hpp index 1bb4993ae0b..1407909c109 100644 --- a/cpp/include/cudf/lists/detail/interleave_columns.hpp +++ b/cpp/include/cudf/lists/detail/interleave_columns.hpp @@ -30,16 +30,17 @@ namespace detail { * @code{.pseudo} * s1 = [{0, 1}, {2, 3, 4}, {5}, {}, {6, 7}] * s2 = [{8}, {9}, {}, {10, 11, 12}, {13, 14, 15, 16}] - * r = interleave_columns(s1, s2) + * r = lists::interleave_columns(s1, s2) * r is now [{0, 1}, {8}, {2, 3, 4}, {9}, {5}, {}, {}, {10, 11, 12}, {6, 7}, {13, 14, 15, 16}] * @endcode * - * @throws cudf::logic_error if columns of the input table are not lists columns. + * @throws cudf::logic_error if any column of the input table is not a lists columns. + * @throws cudf::logic_error if any lists column contains nested type. + * @throws cudf::logic_error if all lists columns do not have the same entry type. * - * @param input Table containing columns to interleave. + * @param input Table containing lists columns to interleave. * @param stream CUDA stream used for device memory operations and kernel launches. - * @param mr Device memory resource used to allocate the returned column's device - * memory. + * @param mr Device memory resource used to allocate the returned column's device memory. * @return The interleaved columns as a single column. */ std::unique_ptr interleave_columns( From dcf95ef7ff2fa0ce7704a4969bc979c82ac9e8f6 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 21 Apr 2021 12:26:15 -0600 Subject: [PATCH 16/36] Add a function to compute the offsets of the output lists column --- cpp/src/lists/interleave_columns.cu | 48 ++++++++++++++++++++++++++--- 1 file changed, 44 insertions(+), 4 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index bc0466aa713..3dac11e89d9 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -20,6 +20,7 @@ #include #include #include +#include #include #include @@ -32,6 +33,46 @@ namespace cudf { namespace lists { namespace detail { namespace { +/** + * @brief Generate list offsets for the output lists column from the table_view of the input lists + * columns. + */ +std::unique_ptr generate_list_offsets(table_view const& lists_columns, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + auto const num_cols = lists_columns.num_columns(); + auto const output_size = lists_columns.num_rows() * num_cols + 1; + + static_assert(sizeof(offset_type) == sizeof(int32_t)); + static_assert(sizeof(size_type) == sizeof(int32_t)); + auto list_offsets = make_numeric_column( + data_type{type_id::INT32}, output_size, mask_state::UNALLOCATED, stream, mr); + auto const d_out_offsets = list_offsets->mutable_view().begin(); + auto const table_dv_ptr = table_device_view::create(lists_columns); + + // Compute list sizes + thrust::transform(rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(output_size - 1), + d_out_offsets, + [num_cols, table_dv = *table_dv_ptr] __device__(size_type const idx) { + auto const col_id = idx % num_cols; + auto const list_id = idx / num_cols; + auto const d_column = table_dv.column(col_id); + if (d_column.is_null(list_id)) { return size_type{0}; } + auto const d_offsets = + d_column.child(lists_column_view::offsets_column_index).data() + + d_column.offset(); + return d_offsets[list_id + 1] - d_offsets[list_id]; + }); + + // Compute offsets from sizes + thrust::exclusive_scan( + rmm::exec_policy(stream), d_out_offsets, d_out_offsets + output_size, d_out_offsets); + + return list_offsets; +} } // anonymous namespace @@ -64,10 +105,9 @@ std::unique_ptr interleave_columns(table_view const& lists_columns, auto const num_rows = lists_columns.num_rows(); if (num_rows == 0) { return cudf::empty_like(lists_columns.column(0)); } - static_assert(sizeof(offset_type) == sizeof(int32_t)); - auto lists_offsets = make_numeric_column( - data_type{type_id::INT32}, num_rows + 1, mask_state::UNALLOCATED, stream, mr); - auto const output_offsets_ptr = lists_offsets->mutable_view().begin(); + auto const list_offsets = generate_list_offsets(lists_columns, stream, mr); + + type_dispatcher(); auto lists_entries = std::make_unique(); From 64dad791b82e43fb66f2a283c7b5559247f18020 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 21 Apr 2021 15:00:10 -0600 Subject: [PATCH 17/36] Fix doxygen again --- cpp/include/cudf/lists/detail/interleave_columns.hpp | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/cpp/include/cudf/lists/detail/interleave_columns.hpp b/cpp/include/cudf/lists/detail/interleave_columns.hpp index 1407909c109..1e21f0800b4 100644 --- a/cpp/include/cudf/lists/detail/interleave_columns.hpp +++ b/cpp/include/cudf/lists/detail/interleave_columns.hpp @@ -38,14 +38,15 @@ namespace detail { * @throws cudf::logic_error if any lists column contains nested type. * @throws cudf::logic_error if all lists columns do not have the same entry type. * - * @param input Table containing lists columns to interleave. + * @param input Table containing lists columns to interleave. + * @param has_null_mask A boolean flag indicating that the input columns have null mask. * @param stream CUDA stream used for device memory operations and kernel launches. - * @param mr Device memory resource used to allocate the returned column's device memory. - * @return The interleaved columns as a single column. + * @param mr Device memory resource used to allocate the returned column's device memory. + * @return The interleaved columns as a single column. */ std::unique_ptr interleave_columns( - table_view const& lists_columns, - bool create_mask, + table_view const& input, + bool has_null_mask, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); From b5dd17c5f2210ef29702971557921b1c85c2f91b Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 21 Apr 2021 15:13:21 -0600 Subject: [PATCH 18/36] Finish skeleton for `lists::detail::interleave_columns` --- cpp/src/lists/interleave_columns.cu | 180 ++++++++++++++++++++++------ 1 file changed, 145 insertions(+), 35 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 3dac11e89d9..1bd39b86d14 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -37,30 +37,45 @@ namespace { * @brief Generate list offsets for the output lists column from the table_view of the input lists * columns. */ -std::unique_ptr generate_list_offsets(table_view const& lists_columns, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +std::pair, rmm::device_uvector> +generate_list_offsets_and_validities(table_view const& input, + bool has_null_mask, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { - auto const num_cols = lists_columns.num_columns(); - auto const output_size = lists_columns.num_rows() * num_cols + 1; + auto const num_cols = input.num_columns(); + auto const output_size = input.num_rows() * num_cols + 1; + // The output offsets column static_assert(sizeof(offset_type) == sizeof(int32_t)); static_assert(sizeof(size_type) == sizeof(int32_t)); auto list_offsets = make_numeric_column( data_type{type_id::INT32}, output_size, mask_state::UNALLOCATED, stream, mr); auto const d_out_offsets = list_offsets->mutable_view().begin(); - auto const table_dv_ptr = table_device_view::create(lists_columns); + auto const table_dv_ptr = table_device_view::create(input); + + // The array of int8_t to store element validities + auto validities = has_null_mask ? rmm::device_uvector(output_size - 1, stream) + : rmm::device_uvector(0, stream); + thrust::uninitialized_fill( + rmm::exec_policy(stream), validities.begin(), validities.end(), int8_t{1}); // Compute list sizes thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(output_size - 1), d_out_offsets, - [num_cols, table_dv = *table_dv_ptr] __device__(size_type const idx) { + [num_cols, + table_dv = *table_dv_ptr, + d_validities = validities.begin(), + has_null_mask] __device__(size_type const idx) { auto const col_id = idx % num_cols; auto const list_id = idx / num_cols; auto const d_column = table_dv.column(col_id); - if (d_column.is_null(list_id)) { return size_type{0}; } + if (d_column.is_null(list_id)) { + if (has_null_mask) { d_validities[idx] = 0; } + return size_type{0}; + } auto const d_offsets = d_column.child(lists_column_view::offsets_column_index).data() + d_column.offset(); @@ -71,60 +86,155 @@ std::unique_ptr generate_list_offsets(table_view const& lists_columns, thrust::exclusive_scan( rmm::exec_policy(stream), d_out_offsets, d_out_offsets + output_size, d_out_offsets); - return list_offsets; + return {std::move(list_offsets), std::move(validities)}; } +/** + * @brief Struct used in type_dispatcher to copy entries to the output lists column + */ +struct copy_list_entries_fn { + template + std::enable_if_t and not cudf::is_fixed_width(), + std::unique_ptr> + operator()(table_view const& input, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) const + { + // Currently, only support string_view and fixed-width types + CUDF_FAIL("Called `copy_list_entries_fn()` on non-supported types."); + } + + template + std::enable_if_t(), std::unique_ptr> operator()( + table_view const& input, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) const noexcept + { + return nullptr; + } + + template + std::enable_if_t, std::unique_ptr> operator()( + table_view const& input, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) const noexcept + { + return nullptr; + } +}; + +/** + * @brief Concatenate strings functor. + * + * This will concatenate the strings within each list element of the given string lists column + * and apply the separator. The null-replacement string scalar `string_narep_dv` (if valid) is + * used in place of any null string. + */ +struct concat_strings_fn { + column_device_view const lists_dv; + offset_type const* const list_offsets; + column_device_view const strings_dv; + string_scalar_device_view const string_narep_dv; + string_view const separator; + + offset_type* d_offsets{nullptr}; + + // If `d_chars == nullptr`: only compute sizes of the output strings. + // If `d_chars != nullptr`: only concatenate strings. + char* d_chars{nullptr}; + + __device__ void operator()(size_type idx) + { + if (lists_dv.is_null(idx)) { + if (!d_chars) { d_offsets[idx] = size_type{-1}; } // negative size means null string + return; + } + + // If the string offsets have been computed and the row `idx` is known to be a null string + if (d_chars && d_offsets[idx] == d_offsets[idx + 1]) { return; } + + auto const separator_size = separator.size_bytes(); + auto size_bytes = size_type{0}; + bool written = false; + char* output_ptr = d_chars ? d_chars + d_offsets[idx] : nullptr; + + for (size_type str_idx = list_offsets[idx], idx_end = list_offsets[idx + 1]; str_idx < idx_end; + ++str_idx) { + if (strings_dv.is_null(str_idx) && !string_narep_dv.is_valid()) { + if (!d_chars) { d_offsets[idx] = size_type{-1}; } // negative size means null string + return; // early termination: the entire list of strings will result in a null string + } + auto const d_str = strings_dv.is_null(str_idx) ? string_narep_dv.value() + : strings_dv.element(str_idx); + size_bytes += separator_size + d_str.size_bytes(); + if (output_ptr) { + // Separator is inserted only in between strings + if (written) { output_ptr = detail::copy_string(output_ptr, separator); } + output_ptr = detail::copy_string(output_ptr, d_str); + written = true; + } + } + + // Separator is inserted only in between strings + if (!d_chars) { d_offsets[idx] = static_cast(size_bytes - separator_size); } + } +}; + } // anonymous namespace /** * @copydoc cudf::lists::detail::interleave_columns * */ -std::unique_ptr interleave_columns(table_view const& lists_columns, - bool create_mask, +std::unique_ptr interleave_columns(table_view const& input, + bool has_null_mask, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - auto const first_col_entry_type_id = - lists_column_view(*lists_columns.begin()).child().type().id(); - for (auto const& col : lists_columns) { + auto const entry_type = lists_column_view(*input.begin()).child().type(); + for (auto const& col : input) { CUDF_EXPECTS(col.type().id() == type_id::LIST, "All columns of the input table must be of lists column type."); auto const child_col = lists_column_view(col).child(); CUDF_EXPECTS(not cudf::is_nested(child_col.type()), "Nested types are not supported."); - CUDF_EXPECTS(first_col_entry_type_id == child_col.type().id(), + CUDF_EXPECTS(entry_type == child_col.type(), "The types of entries in the input columns must be the same."); } - // Single column returns a copy - if (lists_columns.num_columns() == 1) { - return std::make_unique(*(lists_columns.begin()), stream, mr); - } + if (input.num_columns() == 1) { return std::make_unique(*(input.begin()), stream, mr); } + if (input.num_rows() == 0) { return cudf::empty_like(input.column(0)); } - auto const num_rows = lists_columns.num_rows(); - if (num_rows == 0) { return cudf::empty_like(lists_columns.column(0)); } + // Generate offsets of the output lists column + auto [list_offsets, list_validities] = + generate_list_offsets_and_validities(input, has_null_mask, stream, mr); - auto const list_offsets = generate_list_offsets(lists_columns, stream, mr); + // Copy entries from the input lists columns to the output lists column - this needed to be + // specialized for different types + auto list_entries = type_dispatcher(entry_type, copy_list_entries_fn{}, input, stream, mr); - type_dispatcher(); + auto const num_output_rows = input.num_rows() * input.num_columns(); - auto lists_entries = std::make_unique(); + if (not has_null_mask) { + return make_lists_column(num_output_rows, + std::move(list_offsets), + std::move(list_entries), + 0, + rmm::device_buffer{}, + stream, + mr); + } - static constexpr auto invalid_size = std::numeric_limits::lowest(); - auto const count_it = thrust::make_counting_iterator(0); - auto [null_mask, null_count] = cudf::detail::valid_if( - count_it, - count_it + num_rows, - [str_sizes = output_offsets_ptr + 1] __device__(size_type idx) { - return str_sizes[idx] != invalid_size; - }, + auto [null_mask, null_count] = cudf::detail::valid_if( + list_validities.begin(), + list_validities.end(), + [] __device__(auto const valid) { return valid; }, stream, mr); - return make_lists_column(num_rows, - std::move(lists_offsets), - std::move(lists_entries), + return make_lists_column(num_output_rows, + std::move(list_offsets), + std::move(list_entries), null_count, null_count ? std::move(null_mask) : rmm::device_buffer{}, stream, From 7908f6465121fb3dfee982790e4779527f050f9d Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 21 Apr 2021 15:26:38 -0600 Subject: [PATCH 19/36] Refactor --- cpp/src/lists/interleave_columns.cu | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 1bd39b86d14..6b42da2f78e 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -43,19 +43,20 @@ generate_list_offsets_and_validities(table_view const& input, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - auto const num_cols = input.num_columns(); - auto const output_size = input.num_rows() * num_cols + 1; + auto const num_cols = input.num_columns(); + auto const num_rows = input.num_columns(); + auto const num_output_rows = num_rows * num_cols; // The output offsets column static_assert(sizeof(offset_type) == sizeof(int32_t)); static_assert(sizeof(size_type) == sizeof(int32_t)); auto list_offsets = make_numeric_column( - data_type{type_id::INT32}, output_size, mask_state::UNALLOCATED, stream, mr); + data_type{type_id::INT32}, num_output_rows + 1, mask_state::UNALLOCATED, stream, mr); auto const d_out_offsets = list_offsets->mutable_view().begin(); auto const table_dv_ptr = table_device_view::create(input); // The array of int8_t to store element validities - auto validities = has_null_mask ? rmm::device_uvector(output_size - 1, stream) + auto validities = has_null_mask ? rmm::device_uvector(num_output_rows, stream) : rmm::device_uvector(0, stream); thrust::uninitialized_fill( rmm::exec_policy(stream), validities.begin(), validities.end(), int8_t{1}); @@ -63,7 +64,7 @@ generate_list_offsets_and_validities(table_view const& input, // Compute list sizes thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), - thrust::make_counting_iterator(output_size - 1), + thrust::make_counting_iterator(num_output_rows), d_out_offsets, [num_cols, table_dv = *table_dv_ptr, @@ -84,7 +85,7 @@ generate_list_offsets_and_validities(table_view const& input, // Compute offsets from sizes thrust::exclusive_scan( - rmm::exec_policy(stream), d_out_offsets, d_out_offsets + output_size, d_out_offsets); + rmm::exec_policy(stream), d_out_offsets, d_out_offsets + num_output_rows + 1, d_out_offsets); return {std::move(list_offsets), std::move(validities)}; } From 8c4b458de6288dcaf659fd667cf43da0a6f1eff0 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Thu, 22 Apr 2021 09:12:35 -0600 Subject: [PATCH 20/36] Implementation for fixed with --- cpp/src/lists/interleave_columns.cu | 190 +++++++++++++++------------- 1 file changed, 100 insertions(+), 90 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 6b42da2f78e..6a9bd5bf6be 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -16,6 +16,7 @@ #include #include +#include #include #include #include @@ -34,8 +35,8 @@ namespace lists { namespace detail { namespace { /** - * @brief Generate list offsets for the output lists column from the table_view of the input lists - * columns. + * @brief Generate list offsets and list validities for the output lists column from the table_view + * of the input lists columns. */ std::pair, rmm::device_uvector> generate_list_offsets_and_validities(table_view const& input, @@ -43,49 +44,49 @@ generate_list_offsets_and_validities(table_view const& input, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - auto const num_cols = input.num_columns(); - auto const num_rows = input.num_columns(); - auto const num_output_rows = num_rows * num_cols; + auto const num_cols = input.num_columns(); + auto const num_rows = input.num_rows(); + auto const num_lists = num_rows * num_cols; // The output offsets column static_assert(sizeof(offset_type) == sizeof(int32_t)); static_assert(sizeof(size_type) == sizeof(int32_t)); auto list_offsets = make_numeric_column( - data_type{type_id::INT32}, num_output_rows + 1, mask_state::UNALLOCATED, stream, mr); + data_type{type_id::INT32}, num_lists + 1, mask_state::UNALLOCATED, stream, mr); auto const d_out_offsets = list_offsets->mutable_view().begin(); auto const table_dv_ptr = table_device_view::create(input); // The array of int8_t to store element validities - auto validities = has_null_mask ? rmm::device_uvector(num_output_rows, stream) + auto validities = has_null_mask ? rmm::device_uvector(num_lists, stream) : rmm::device_uvector(0, stream); - thrust::uninitialized_fill( - rmm::exec_policy(stream), validities.begin(), validities.end(), int8_t{1}); // Compute list sizes thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), - thrust::make_counting_iterator(num_output_rows), + thrust::make_counting_iterator(num_lists), d_out_offsets, [num_cols, table_dv = *table_dv_ptr, d_validities = validities.begin(), - has_null_mask] __device__(size_type const idx) { - auto const col_id = idx % num_cols; - auto const list_id = idx / num_cols; - auto const d_column = table_dv.column(col_id); - if (d_column.is_null(list_id)) { - if (has_null_mask) { d_validities[idx] = 0; } - return size_type{0}; + has_null_mask] __device__(size_type const dst_list_id) { + auto const src_col_id = dst_list_id % num_cols; + auto const src_list_id = dst_list_id / num_cols; + auto const& src_col = table_dv.column(src_col_id); + auto const is_valid = src_col.is_valid(src_list_id); + + if (has_null_mask) { + d_validities[dst_list_id] = static_cast(is_valid); } + if (not is_valid) { return size_type{0}; } auto const d_offsets = - d_column.child(lists_column_view::offsets_column_index).data() + - d_column.offset(); - return d_offsets[list_id + 1] - d_offsets[list_id]; + src_col.child(lists_column_view::offsets_column_index).data() + + src_col.offset(); + return d_offsets[src_list_id + 1] - d_offsets[src_list_id]; }); // Compute offsets from sizes thrust::exclusive_scan( - rmm::exec_policy(stream), d_out_offsets, d_out_offsets + num_output_rows + 1, d_out_offsets); + rmm::exec_policy(stream), d_out_offsets, d_out_offsets + num_lists + 1, d_out_offsets); return {std::move(list_offsets), std::move(validities)}; } @@ -95,19 +96,10 @@ generate_list_offsets_and_validities(table_view const& input, */ struct copy_list_entries_fn { template - std::enable_if_t and not cudf::is_fixed_width(), - std::unique_ptr> - operator()(table_view const& input, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) const - { - // Currently, only support string_view and fixed-width types - CUDF_FAIL("Called `copy_list_entries_fn()` on non-supported types."); - } - - template - std::enable_if_t(), std::unique_ptr> operator()( + std::enable_if_t, std::unique_ptr> operator()( table_view const& input, + column_view const& list_offsets, + bool has_null_mask, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) const noexcept { @@ -115,69 +107,87 @@ struct copy_list_entries_fn { } template - std::enable_if_t, std::unique_ptr> operator()( + std::enable_if_t(), std::unique_ptr> operator()( table_view const& input, + column_view const& list_offsets, + bool has_null_mask, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) const noexcept { - return nullptr; - } -}; + auto const child_col = lists_column_view(*input.begin()).child(); + auto const num_cols = input.num_columns(); + auto const num_rows = input.num_rows(); + auto const num_lists = num_rows * num_cols; + auto const output_size = cudf::detail::get_value(list_offsets, num_lists, stream); -/** - * @brief Concatenate strings functor. - * - * This will concatenate the strings within each list element of the given string lists column - * and apply the separator. The null-replacement string scalar `string_narep_dv` (if valid) is - * used in place of any null string. - */ -struct concat_strings_fn { - column_device_view const lists_dv; - offset_type const* const list_offsets; - column_device_view const strings_dv; - string_scalar_device_view const string_narep_dv; - string_view const separator; + std::vector child_views, offsets_views; + for (auto const& col : input) { + auto const list_col = lists_column_view(col); + child_views.emplace_back(list_col.child()); + offsets_views.emplace_back(list_col.offsets()); + } + auto const child_cols_ptr = table_device_view::create(table_view{child_views}); + auto const offsets_cols_ptr = table_device_view::create(table_view{offsets_views}); - offset_type* d_offsets{nullptr}; + auto output = allocate_like(child_col, output_size, mask_allocation_policy::NEVER, stream, mr); + auto output_dv_ptr = mutable_column_device_view::create(*output); - // If `d_chars == nullptr`: only compute sizes of the output strings. - // If `d_chars != nullptr`: only concatenate strings. - char* d_chars{nullptr}; + // The array of int8_t to store element validities + auto validities = has_null_mask ? rmm::device_uvector(output_size, stream) + : rmm::device_uvector(0, stream); - __device__ void operator()(size_type idx) - { - if (lists_dv.is_null(idx)) { - if (!d_chars) { d_offsets[idx] = size_type{-1}; } // negative size means null string - return; - } + thrust::for_each_n(rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + num_lists, + [num_cols, + child_cols = *child_cols_ptr, + offsets_cols = *offsets_cols_ptr, + d_validities = validities.begin(), + dst_offsets = list_offsets.begin(), + d_output = output_dv_ptr->begin(), + has_null_mask] __device__(size_type const dst_list_id) { + auto const src_col_id = dst_list_id % num_cols; + auto const src_list_id = dst_list_id / num_cols; + auto const& src_col = child_cols.column(src_col_id); + auto const& src_offsets = offsets_cols.column(src_col_id); - // If the string offsets have been computed and the row `idx` is known to be a null string - if (d_chars && d_offsets[idx] == d_offsets[idx + 1]) { return; } - - auto const separator_size = separator.size_bytes(); - auto size_bytes = size_type{0}; - bool written = false; - char* output_ptr = d_chars ? d_chars + d_offsets[idx] : nullptr; - - for (size_type str_idx = list_offsets[idx], idx_end = list_offsets[idx + 1]; str_idx < idx_end; - ++str_idx) { - if (strings_dv.is_null(str_idx) && !string_narep_dv.is_valid()) { - if (!d_chars) { d_offsets[idx] = size_type{-1}; } // negative size means null string - return; // early termination: the entire list of strings will result in a null string - } - auto const d_str = strings_dv.is_null(str_idx) ? string_narep_dv.value() - : strings_dv.element(str_idx); - size_bytes += separator_size + d_str.size_bytes(); - if (output_ptr) { - // Separator is inserted only in between strings - if (written) { output_ptr = detail::copy_string(output_ptr, separator); } - output_ptr = detail::copy_string(output_ptr, d_str); - written = true; - } + auto write_idx = dst_offsets[dst_list_id]; + for (auto read_idx = src_offsets.element(src_list_id), + idx_end = src_offsets.element(src_list_id + 1); + read_idx < idx_end; + ++read_idx, ++write_idx) { + auto const is_valid = src_col.is_valid(read_idx); + if (has_null_mask) { + d_validities[dst_list_id] = static_cast(is_valid); + } + d_output[write_idx] = is_valid ? src_col.element(read_idx) : T{}; + } + }); + + if (has_null_mask) { + auto [null_mask, null_count] = cudf::detail::valid_if( + validities.begin(), + validities.end(), + [] __device__(auto const valid) { return valid; }, + stream, + mr); + if (null_count > 0) { output->set_null_mask(null_mask, null_count); } } - // Separator is inserted only in between strings - if (!d_chars) { d_offsets[idx] = static_cast(size_bytes - separator_size); } + return output; + } + + template + std::enable_if_t and not cudf::is_fixed_width(), + std::unique_ptr> + operator()(table_view const&, + column_view const&, + bool, + rmm::cuda_stream_view, + rmm::mr::device_memory_resource*) const + { + // Currently, only support string_view and fixed-width types + CUDF_FAIL("Called `copy_list_entries_fn()` on non-supported types."); } }; @@ -212,12 +222,12 @@ std::unique_ptr interleave_columns(table_view const& input, // Copy entries from the input lists columns to the output lists column - this needed to be // specialized for different types - auto list_entries = type_dispatcher(entry_type, copy_list_entries_fn{}, input, stream, mr); - - auto const num_output_rows = input.num_rows() * input.num_columns(); + auto const output_size = input.num_rows() * input.num_columns(); + auto list_entries = type_dispatcher( + entry_type, copy_list_entries_fn{}, input, list_offsets->view(), has_null_mask, stream, mr); if (not has_null_mask) { - return make_lists_column(num_output_rows, + return make_lists_column(output_size, std::move(list_offsets), std::move(list_entries), 0, @@ -233,7 +243,7 @@ std::unique_ptr interleave_columns(table_view const& input, stream, mr); - return make_lists_column(num_output_rows, + return make_lists_column(output_size, std::move(list_offsets), std::move(list_entries), null_count, From 45ed175428a389847912b04ca9070c6a3fa2cbe5 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Thu, 22 Apr 2021 10:57:04 -0600 Subject: [PATCH 21/36] Fix error --- cpp/src/lists/interleave_columns.cu | 90 ++++++++++++++++------------- 1 file changed, 49 insertions(+), 41 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 6a9bd5bf6be..651a183245f 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -30,6 +30,8 @@ #include +#include + namespace cudf { namespace lists { namespace detail { @@ -114,55 +116,56 @@ struct copy_list_entries_fn { rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) const noexcept { - auto const child_col = lists_column_view(*input.begin()).child(); - auto const num_cols = input.num_columns(); - auto const num_rows = input.num_rows(); - auto const num_lists = num_rows * num_cols; - auto const output_size = cudf::detail::get_value(list_offsets, num_lists, stream); - - std::vector child_views, offsets_views; - for (auto const& col : input) { - auto const list_col = lists_column_view(col); - child_views.emplace_back(list_col.child()); - offsets_views.emplace_back(list_col.offsets()); - } - auto const child_cols_ptr = table_device_view::create(table_view{child_views}); - auto const offsets_cols_ptr = table_device_view::create(table_view{offsets_views}); + auto const child_col = lists_column_view(*input.begin()).child(); + auto const num_cols = input.num_columns(); + auto const num_rows = input.num_rows(); + auto const num_lists = num_rows * num_cols; + auto const output_size = cudf::detail::get_value(list_offsets, num_lists, stream); + auto const table_dv_ptr = table_device_view::create(input); + + printf("Line %d\n", __LINE__); + + printf("Line %d\n", __LINE__); + + printf("Line %d\n", __LINE__); auto output = allocate_like(child_col, output_size, mask_allocation_policy::NEVER, stream, mr); auto output_dv_ptr = mutable_column_device_view::create(*output); + printf("Line %d\n", __LINE__); + // The array of int8_t to store element validities auto validities = has_null_mask ? rmm::device_uvector(output_size, stream) : rmm::device_uvector(0, stream); - thrust::for_each_n(rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - num_lists, - [num_cols, - child_cols = *child_cols_ptr, - offsets_cols = *offsets_cols_ptr, - d_validities = validities.begin(), - dst_offsets = list_offsets.begin(), - d_output = output_dv_ptr->begin(), - has_null_mask] __device__(size_type const dst_list_id) { - auto const src_col_id = dst_list_id % num_cols; - auto const src_list_id = dst_list_id / num_cols; - auto const& src_col = child_cols.column(src_col_id); - auto const& src_offsets = offsets_cols.column(src_col_id); - - auto write_idx = dst_offsets[dst_list_id]; - for (auto read_idx = src_offsets.element(src_list_id), - idx_end = src_offsets.element(src_list_id + 1); - read_idx < idx_end; - ++read_idx, ++write_idx) { - auto const is_valid = src_col.is_valid(read_idx); - if (has_null_mask) { - d_validities[dst_list_id] = static_cast(is_valid); - } - d_output[write_idx] = is_valid ? src_col.element(read_idx) : T{}; - } - }); + thrust::for_each_n( + rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + num_lists, + [num_cols, + table_dv = *table_dv_ptr, + d_validities = validities.begin(), + dst_offsets = list_offsets.begin(), + d_output = output_dv_ptr->begin(), + has_null_mask] __device__(size_type const dst_list_id) { + auto const src_col_id = dst_list_id % num_cols; + auto const src_list_id = dst_list_id / num_cols; + auto const& src_lists_col = table_dv.column(src_col_id); + auto const& src_child = src_lists_col.child(lists_column_view::child_column_index); + auto const src_child_offsets = + src_lists_col.child(lists_column_view::offsets_column_index).data() + + src_lists_col.offset(); + + auto write_idx = dst_offsets[dst_list_id]; + for (auto read_idx = src_child_offsets[src_list_id], + idx_end = src_child_offsets[src_list_id + 1]; + read_idx < idx_end; + ++read_idx, ++write_idx) { + auto const is_valid = src_child.is_valid(read_idx); + if (has_null_mask) { d_validities[dst_list_id] = static_cast(is_valid); } + d_output[write_idx] = is_valid ? src_child.element(read_idx) : T{}; + } + }); if (has_null_mask) { auto [null_mask, null_count] = cudf::detail::valid_if( @@ -220,12 +223,16 @@ std::unique_ptr interleave_columns(table_view const& input, auto [list_offsets, list_validities] = generate_list_offsets_and_validities(input, has_null_mask, stream, mr); + printf("Line %d\n", __LINE__); + // Copy entries from the input lists columns to the output lists column - this needed to be // specialized for different types auto const output_size = input.num_rows() * input.num_columns(); auto list_entries = type_dispatcher( entry_type, copy_list_entries_fn{}, input, list_offsets->view(), has_null_mask, stream, mr); + printf("Line %d\n", __LINE__); + if (not has_null_mask) { return make_lists_column(output_size, std::move(list_offsets), @@ -242,6 +249,7 @@ std::unique_ptr interleave_columns(table_view const& input, [] __device__(auto const valid) { return valid; }, stream, mr); + printf("Line %d\n", __LINE__); return make_lists_column(output_size, std::move(list_offsets), From 9ac01437d643a9ac30784aa26e4154f2410da8b8 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Thu, 22 Apr 2021 15:56:48 -0600 Subject: [PATCH 22/36] Draft --- cpp/src/lists/interleave_columns.cu | 242 +++++++++++++++++++++++----- 1 file changed, 205 insertions(+), 37 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 651a183245f..7d3c1727732 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -14,6 +14,8 @@ * limitations under the License. */ +#include + #include #include #include @@ -28,6 +30,7 @@ #include #include +#include #include #include @@ -75,7 +78,6 @@ generate_list_offsets_and_validities(table_view const& input, auto const src_list_id = dst_list_id / num_cols; auto const& src_col = table_dv.column(src_col_id); auto const is_valid = src_col.is_valid(src_list_id); - if (has_null_mask) { d_validities[dst_list_id] = static_cast(is_valid); } @@ -93,6 +95,154 @@ generate_list_offsets_and_validities(table_view const& input, return {std::move(list_offsets), std::move(validities)}; } +/** + * @brief Creates child offsets, chars columns and null mask, null count of a strings column by + * applying the template function that can be used for computing the output size of each string as + * well as create the output. + * + * @tparam SizeAndExecuteFunction Function must accept an index and return a size. + * It must have members `d_offsets`, `d_chars`, and `d_validities` which are set to memory + * containing the offsets column, chars column and string validities during write. + * + * @param size_and_exec_fn This is called twice. Once for the output size of each string, which is + * written into the `d_offsets` array. After that, `d_chars` is set and this + * is called again to fill in the chars memory. The `d_validities` array may + * be modified to set the value `0` for the corresponding rows that contain + * null string elements. + * @param strings_count Number of strings. + * @param mr Device memory resource used to allocate the returned columns' device memory. + * @param stream CUDA stream used for device memory operations and kernel launches. + * @return offsets child column, chars child column, null_mask, and null_count for a strings column. + */ +template +std::tuple, std::unique_ptr, rmm::device_buffer, size_type> +make_strings_children_with_null_mask( + SizeAndExecuteFunction size_and_exec_fn, + size_type strings_count, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +{ + auto offsets_column = make_numeric_column( + data_type{type_id::INT32}, strings_count + 1, mask_state::UNALLOCATED, stream, mr); + auto offsets_view = offsets_column->mutable_view(); + auto d_offsets = offsets_view.template data(); + size_and_exec_fn.d_offsets = d_offsets; + + auto validities = rmm::device_uvector(strings_count, stream); + size_and_exec_fn.d_validities = validities.begin(); + + // This is called twice: once for offsets and validities, and once for chars + auto for_each_fn = [strings_count, stream](SizeAndExecuteFunction& size_and_exec_fn) { + thrust::for_each_n(rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + strings_count, + size_and_exec_fn); + }; + + // Compute the string sizes (storing in `d_offsets`) and string validities + for_each_fn(size_and_exec_fn); + + // Compute the offsets from string sizes + thrust::exclusive_scan( + rmm::exec_policy(stream), d_offsets, d_offsets + strings_count + 1, d_offsets); + + // Now build the chars column + auto const bytes = cudf::detail::get_value(offsets_view, strings_count, stream); + auto chars_column = + cudf::strings::detail::create_chars_child_column(strings_count, 0, bytes, stream, mr); + + // Execute the function fn again to fill the chars column. + // Note that if the output chars column has zero size, the function fn should not be called to + // avoid accidentally overwriting the offsets. + if (bytes > 0) { + size_and_exec_fn.d_chars = chars_column->mutable_view().template data(); + for_each_fn(size_and_exec_fn); + } + + // Finally compute null mask and null count from the validities array + auto [null_mask, null_count] = cudf::detail::valid_if( + validities.begin(), + validities.end(), + [] __device__(auto const valid) { return valid; }, + stream, + mr); + + return std::make_tuple(std::move(offsets_column), + std::move(chars_column), + null_count > 0 ? std::move(null_mask) : rmm::device_buffer{}, + null_count); +} + +/** + * @brief Compute string sizes, string validities, and concatenate strings functor. + * + * This functor is executed twice. In the first pass, the sizes and validities of the output strings + * will be computed. In the second pass, this will concatenate the strings within each list element + * of the given lists column and apply the separator. The null-replacement string scalar + * `string_narep_dv` (if valid) is used in place of any null string. + * + * @tparam Functor The functor which can check for validity of the input list at a given list index + * as well as access to the separator corresponding to the list index. + */ +struct compute_size_and_copy_fn { + table_device_view const table_dv; + + // Store offsets of the lists in the output lists column + offset_type const* const dst_list_offsets; + + // Store offsets of the strings in the chars column + offset_type* d_offsets{nullptr}; + + // If d_chars == nullptr: only compute sizes and validities of the output strings. + // If d_chars != nullptr: only concatenate strings. + char* d_chars{nullptr}; + + // We need to set `1` or `0` for the validities of the strings in the child column. + int8_t* d_validities{nullptr}; + + bool has_null_mask; + + __device__ void operator()(size_type const dst_list_id) + { + auto const num_cols = table_dv.num_columns(); + auto const src_col_id = dst_list_id % num_cols; + auto const src_list_id = dst_list_id / num_cols; + auto const& src_lists_col = table_dv.column(src_col_id); + if (has_null_mask and src_lists_col.is_null(src_list_id)) { return; } + + auto const src_list_offsets = + src_lists_col.child(lists_column_view::offsets_column_index).data() + + src_lists_col.offset(); + auto const& src_child = src_lists_col.child(lists_column_view::child_column_index); + auto const src_child_offsets = + src_child.child(strings_column_view::offsets_column_index).data() + + src_child.offset(); + + size_type write_idx = dst_list_offsets[dst_list_id]; + if (not d_chars) { // just compute sizes of strings within a list + for (auto read_idx = src_list_offsets[src_list_id], + end_idx = src_list_offsets[src_list_id + 1]; + read_idx < end_idx; + ++read_idx, ++write_idx) { + auto const is_valid = src_child.is_valid(read_idx); + if (has_null_mask) { d_validities[write_idx] = static_cast(is_valid); } + d_offsets[write_idx] = + is_valid ? src_child_offsets[read_idx + 1] - src_child_offsets[read_idx] : 0; + } + } else { // just copy the entire list of strings + auto const start_idx = src_child_offsets[src_list_offsets[src_list_id]]; + auto const end_idx = src_child_offsets[src_list_offsets[src_list_id + 1] - 1]; + if (start_idx < end_idx) { + auto const input_ptr = + src_child.child(strings_column_view::chars_column_index).data() + + src_child.offset() + start_idx; + auto const output_ptr = d_chars + d_offsets[write_idx]; + std::memcpy(output_ptr, input_ptr, end_idx - start_idx); + } + } + } +}; + /** * @brief Struct used in type_dispatcher to copy entries to the output lists column */ @@ -101,17 +251,41 @@ struct copy_list_entries_fn { std::enable_if_t, std::unique_ptr> operator()( table_view const& input, column_view const& list_offsets, + size_type num_strings, bool has_null_mask, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) const noexcept { - return nullptr; + auto const table_dv_ptr = table_device_view::create(input); + auto const comp_fn = compute_size_and_copy_fn{*table_dv_ptr, list_offsets.begin()}; + if (not has_null_mask) { + auto [offsets_column, chars_column] = + cudf::strings::detail::make_strings_children(comp_fn, num_strings, 0, stream, mr); + return make_strings_column(num_strings, + std::move(offsets_column), + std::move(chars_column), + 0, + rmm::device_buffer{}, + stream, + mr); + } + + auto [offsets_column, chars_column, null_mask, null_count] = + make_strings_children_with_null_mask(comp_fn, num_strings, stream, mr); + return make_strings_column(num_strings, + std::move(offsets_column), + std::move(chars_column), + null_count, + std::move(null_mask), + stream, + mr); } template std::enable_if_t(), std::unique_ptr> operator()( table_view const& input, column_view const& list_offsets, + size_type num_output_entries, bool has_null_mask, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) const noexcept @@ -120,22 +294,14 @@ struct copy_list_entries_fn { auto const num_cols = input.num_columns(); auto const num_rows = input.num_rows(); auto const num_lists = num_rows * num_cols; - auto const output_size = cudf::detail::get_value(list_offsets, num_lists, stream); auto const table_dv_ptr = table_device_view::create(input); - printf("Line %d\n", __LINE__); - - printf("Line %d\n", __LINE__); - - printf("Line %d\n", __LINE__); - - auto output = allocate_like(child_col, output_size, mask_allocation_policy::NEVER, stream, mr); + auto output = + allocate_like(child_col, num_output_entries, mask_allocation_policy::NEVER, stream, mr); auto output_dv_ptr = mutable_column_device_view::create(*output); - printf("Line %d\n", __LINE__); - // The array of int8_t to store element validities - auto validities = has_null_mask ? rmm::device_uvector(output_size, stream) + auto validities = has_null_mask ? rmm::device_uvector(num_output_entries, stream) : rmm::device_uvector(0, stream); thrust::for_each_n( @@ -143,27 +309,26 @@ struct copy_list_entries_fn { thrust::make_counting_iterator(0), num_lists, [num_cols, - table_dv = *table_dv_ptr, - d_validities = validities.begin(), - dst_offsets = list_offsets.begin(), - d_output = output_dv_ptr->begin(), + table_dv = *table_dv_ptr, + out_validities = validities.begin(), + dst_list_offsets = list_offsets.begin(), + out_entries = output_dv_ptr->begin(), has_null_mask] __device__(size_type const dst_list_id) { auto const src_col_id = dst_list_id % num_cols; auto const src_list_id = dst_list_id / num_cols; auto const& src_lists_col = table_dv.column(src_col_id); - auto const& src_child = src_lists_col.child(lists_column_view::child_column_index); - auto const src_child_offsets = + auto const src_list_offsets = src_lists_col.child(lists_column_view::offsets_column_index).data() + src_lists_col.offset(); - - auto write_idx = dst_offsets[dst_list_id]; - for (auto read_idx = src_child_offsets[src_list_id], - idx_end = src_child_offsets[src_list_id + 1]; - read_idx < idx_end; + auto const& src_child = src_lists_col.child(lists_column_view::child_column_index); + for (auto read_idx = src_list_offsets[src_list_id], + end_idx = src_list_offsets[src_list_id + 1], + write_idx = dst_list_offsets[dst_list_id]; + read_idx < end_idx; ++read_idx, ++write_idx) { auto const is_valid = src_child.is_valid(read_idx); - if (has_null_mask) { d_validities[dst_list_id] = static_cast(is_valid); } - d_output[write_idx] = is_valid ? src_child.element(read_idx) : T{}; + if (has_null_mask) { out_validities[write_idx] = static_cast(is_valid); } + out_entries[write_idx] = is_valid ? src_child.element(read_idx) : T{}; } }); @@ -185,6 +350,7 @@ struct copy_list_entries_fn { std::unique_ptr> operator()(table_view const&, column_view const&, + size_type, bool, rmm::cuda_stream_view, rmm::mr::device_memory_resource*) const @@ -223,18 +389,22 @@ std::unique_ptr interleave_columns(table_view const& input, auto [list_offsets, list_validities] = generate_list_offsets_and_validities(input, has_null_mask, stream, mr); - printf("Line %d\n", __LINE__); - // Copy entries from the input lists columns to the output lists column - this needed to be // specialized for different types - auto const output_size = input.num_rows() * input.num_columns(); - auto list_entries = type_dispatcher( - entry_type, copy_list_entries_fn{}, input, list_offsets->view(), has_null_mask, stream, mr); - - printf("Line %d\n", __LINE__); + auto const num_output_lists = input.num_rows() * input.num_columns(); + auto const num_output_entries = + cudf::detail::get_value(list_offsets->view(), num_output_lists, stream); + auto list_entries = type_dispatcher(entry_type, + copy_list_entries_fn{}, + input, + list_offsets->view(), + num_output_entries, + has_null_mask, + stream, + mr); if (not has_null_mask) { - return make_lists_column(output_size, + return make_lists_column(num_output_lists, std::move(list_offsets), std::move(list_entries), 0, @@ -249,9 +419,7 @@ std::unique_ptr interleave_columns(table_view const& input, [] __device__(auto const valid) { return valid; }, stream, mr); - printf("Line %d\n", __LINE__); - - return make_lists_column(output_size, + return make_lists_column(num_output_lists, std::move(list_offsets), std::move(list_entries), null_count, From a406543b851b74e2eb8950bd9b46bc75552c83e3 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Thu, 22 Apr 2021 20:33:24 -0600 Subject: [PATCH 23/36] Fix a bug in unit tests --- cpp/tests/reshape/interleave_columns_tests.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index 56f29033842..62ec285894f 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -592,10 +592,10 @@ TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputNoNull) TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputWithNulls) { auto const col_original = ListsCol{{ListsCol{{null, 2, 3}, null_at(0)}, - ListsCol{{2, 3}}, /*NULL*/ + ListsCol{2, 3}, /*NULL*/ ListsCol{{3, null, 5, 6}, null_at(1)}, - ListsCol{{5, 6}}, /*NULL*/ - ListsCol{}, /*NULL*/ + ListsCol{5, 6}, /*NULL*/ + ListsCol{}, /*NULL*/ ListsCol{7}, ListsCol{8, 9, 10}}, null_at({1, 3, 4})} From e1c0a5a16d87a9f090b1a13867bc4a170de0903b Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Thu, 22 Apr 2021 20:33:38 -0600 Subject: [PATCH 24/36] working --- cpp/src/lists/interleave_columns.cu | 96 ++++++++++++++++++++--------- 1 file changed, 66 insertions(+), 30 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 7d3c1727732..1c48fe2efde 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -118,7 +118,9 @@ template std::tuple, std::unique_ptr, rmm::device_buffer, size_type> make_strings_children_with_null_mask( SizeAndExecuteFunction size_and_exec_fn, + size_type lists_count, size_type strings_count, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) { @@ -132,20 +134,34 @@ make_strings_children_with_null_mask( size_and_exec_fn.d_validities = validities.begin(); // This is called twice: once for offsets and validities, and once for chars - auto for_each_fn = [strings_count, stream](SizeAndExecuteFunction& size_and_exec_fn) { + auto for_each_fn = [lists_count, stream](SizeAndExecuteFunction& size_and_exec_fn) { thrust::for_each_n(rmm::exec_policy(stream), thrust::make_counting_iterator(0), - strings_count, + lists_count, size_and_exec_fn); }; // Compute the string sizes (storing in `d_offsets`) and string validities for_each_fn(size_and_exec_fn); + printf("Line %d\n", __LINE__); + std::vector v(validities.size()); + CUDA_TRY(cudaMemcpyAsync( + v.data(), validities.data(), validities.size(), cudaMemcpyDefault, stream.value())); + stream.synchronize(); + for (auto x : v) { printf("%d, ", (int)x); } + printf("\n\n"); + + printf("Line %d\n", __LINE__); + cudf::test::print(offsets_view); + // Compute the offsets from string sizes thrust::exclusive_scan( rmm::exec_policy(stream), d_offsets, d_offsets + strings_count + 1, d_offsets); + printf("Line %d\n", __LINE__); + cudf::test::print(offsets_view); + // Now build the chars column auto const bytes = cudf::detail::get_value(offsets_view, strings_count, stream); auto chars_column = @@ -190,6 +206,8 @@ struct compute_size_and_copy_fn { // Store offsets of the lists in the output lists column offset_type const* const dst_list_offsets; + bool const has_null_mask; + // Store offsets of the strings in the chars column offset_type* d_offsets{nullptr}; @@ -200,15 +218,17 @@ struct compute_size_and_copy_fn { // We need to set `1` or `0` for the validities of the strings in the child column. int8_t* d_validities{nullptr}; - bool has_null_mask; - __device__ void operator()(size_type const dst_list_id) { - auto const num_cols = table_dv.num_columns(); - auto const src_col_id = dst_list_id % num_cols; - auto const src_list_id = dst_list_id / num_cols; + auto const num_cols = table_dv.num_columns(); + auto const src_col_id = dst_list_id % num_cols; + auto const src_list_id = dst_list_id / num_cols; + auto const& src_lists_col = table_dv.column(src_col_id); - if (has_null_mask and src_lists_col.is_null(src_list_id)) { return; } + if (has_null_mask and src_lists_col.is_null(src_list_id)) { + printf("return c null\n"); + return; + } auto const src_list_offsets = src_lists_col.child(lists_column_view::offsets_column_index).data() + @@ -226,12 +246,11 @@ struct compute_size_and_copy_fn { ++read_idx, ++write_idx) { auto const is_valid = src_child.is_valid(read_idx); if (has_null_mask) { d_validities[write_idx] = static_cast(is_valid); } - d_offsets[write_idx] = - is_valid ? src_child_offsets[read_idx + 1] - src_child_offsets[read_idx] : 0; + d_offsets[write_idx] = src_child_offsets[read_idx + 1] - src_child_offsets[read_idx]; } } else { // just copy the entire list of strings auto const start_idx = src_child_offsets[src_list_offsets[src_list_id]]; - auto const end_idx = src_child_offsets[src_list_offsets[src_list_id + 1] - 1]; + auto const end_idx = src_child_offsets[src_list_offsets[src_list_id + 1]]; if (start_idx < end_idx) { auto const input_ptr = src_child.child(strings_column_view::chars_column_index).data() + @@ -251,40 +270,56 @@ struct copy_list_entries_fn { std::enable_if_t, std::unique_ptr> operator()( table_view const& input, column_view const& list_offsets, - size_type num_strings, + size_type num_output_lists, + size_type num_output_entries, bool has_null_mask, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) const noexcept { auto const table_dv_ptr = table_device_view::create(input); - auto const comp_fn = compute_size_and_copy_fn{*table_dv_ptr, list_offsets.begin()}; - if (not has_null_mask) { - auto [offsets_column, chars_column] = - cudf::strings::detail::make_strings_children(comp_fn, num_strings, 0, stream, mr); - return make_strings_column(num_strings, + auto const comp_fn = + compute_size_and_copy_fn{*table_dv_ptr, list_offsets.begin(), has_null_mask}; + // if (not has_null_mask) { + // auto [offsets_column, chars_column] = + // cudf::strings::detail::make_strings_children(comp_fn, num_output_lists, 0, stream, + // mr); + // return make_strings_column(num_output_entries, + // std::move(offsets_column), + // std::move(chars_column), + // 0, + // rmm::device_buffer{}, + // stream, + // mr); + // } + + printf("numlist: %d, entri: %d\n", num_output_lists, num_output_entries); + + auto [offsets_column, chars_column, null_mask, null_count] = + make_strings_children_with_null_mask( + comp_fn, num_output_lists, num_output_entries, stream, mr); + if (not has_null_mask) + return make_strings_column(num_output_entries, std::move(offsets_column), std::move(chars_column), 0, rmm::device_buffer{}, stream, mr); - } - - auto [offsets_column, chars_column, null_mask, null_count] = - make_strings_children_with_null_mask(comp_fn, num_strings, stream, mr); - return make_strings_column(num_strings, - std::move(offsets_column), - std::move(chars_column), - null_count, - std::move(null_mask), - stream, - mr); + else + return make_strings_column(num_output_entries, + std::move(offsets_column), + std::move(chars_column), + null_count, + std::move(null_mask), + stream, + mr); } template std::enable_if_t(), std::unique_ptr> operator()( table_view const& input, column_view const& list_offsets, + size_type num_output_lists, size_type num_output_entries, bool has_null_mask, rmm::cuda_stream_view stream, @@ -293,7 +328,6 @@ struct copy_list_entries_fn { auto const child_col = lists_column_view(*input.begin()).child(); auto const num_cols = input.num_columns(); auto const num_rows = input.num_rows(); - auto const num_lists = num_rows * num_cols; auto const table_dv_ptr = table_device_view::create(input); auto output = @@ -307,7 +341,7 @@ struct copy_list_entries_fn { thrust::for_each_n( rmm::exec_policy(stream), thrust::make_counting_iterator(0), - num_lists, + num_output_lists, [num_cols, table_dv = *table_dv_ptr, out_validities = validities.begin(), @@ -351,6 +385,7 @@ struct copy_list_entries_fn { operator()(table_view const&, column_view const&, size_type, + size_type, bool, rmm::cuda_stream_view, rmm::mr::device_memory_resource*) const @@ -398,6 +433,7 @@ std::unique_ptr interleave_columns(table_view const& input, copy_list_entries_fn{}, input, list_offsets->view(), + num_output_lists, num_output_entries, has_null_mask, stream, From ac31b2eb93e070b5d06d3efc85d5fdd7cf6e14de Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Fri, 23 Apr 2021 08:31:02 -0600 Subject: [PATCH 25/36] Complete --- .../cudf/lists/detail/interleave_columns.hpp | 2 +- cpp/src/lists/interleave_columns.cu | 301 ++++++------------ 2 files changed, 99 insertions(+), 204 deletions(-) diff --git a/cpp/include/cudf/lists/detail/interleave_columns.hpp b/cpp/include/cudf/lists/detail/interleave_columns.hpp index 1e21f0800b4..831ca5d885b 100644 --- a/cpp/include/cudf/lists/detail/interleave_columns.hpp +++ b/cpp/include/cudf/lists/detail/interleave_columns.hpp @@ -35,7 +35,7 @@ namespace detail { * @endcode * * @throws cudf::logic_error if any column of the input table is not a lists columns. - * @throws cudf::logic_error if any lists column contains nested type. + * @throws cudf::logic_error if any lists column contains nested typed entry. * @throws cudf::logic_error if all lists columns do not have the same entry type. * * @param input Table containing lists columns to interleave. diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 1c48fe2efde..6abec23383d 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -14,15 +14,12 @@ * limitations under the License. */ -#include - #include #include #include -#include -#include #include -#include +#include +#include #include #include @@ -30,11 +27,8 @@ #include #include -#include #include -#include - namespace cudf { namespace lists { namespace detail { @@ -49,166 +43,67 @@ generate_list_offsets_and_validities(table_view const& input, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - auto const num_cols = input.num_columns(); - auto const num_rows = input.num_rows(); - auto const num_lists = num_rows * num_cols; + auto const num_cols = input.num_columns(); + auto const num_rows = input.num_rows(); + auto const num_output_lists = num_rows * num_cols; + auto const table_dv_ptr = table_device_view::create(input); - // The output offsets column + // The output offsets column. static_assert(sizeof(offset_type) == sizeof(int32_t)); static_assert(sizeof(size_type) == sizeof(int32_t)); auto list_offsets = make_numeric_column( - data_type{type_id::INT32}, num_lists + 1, mask_state::UNALLOCATED, stream, mr); - auto const d_out_offsets = list_offsets->mutable_view().begin(); - auto const table_dv_ptr = table_device_view::create(input); - - // The array of int8_t to store element validities - auto validities = has_null_mask ? rmm::device_uvector(num_lists, stream) - : rmm::device_uvector(0, stream); - - // Compute list sizes - thrust::transform(rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(num_lists), - d_out_offsets, - [num_cols, - table_dv = *table_dv_ptr, - d_validities = validities.begin(), - has_null_mask] __device__(size_type const dst_list_id) { - auto const src_col_id = dst_list_id % num_cols; - auto const src_list_id = dst_list_id / num_cols; - auto const& src_col = table_dv.column(src_col_id); - auto const is_valid = src_col.is_valid(src_list_id); - if (has_null_mask) { - d_validities[dst_list_id] = static_cast(is_valid); - } - if (not is_valid) { return size_type{0}; } - auto const d_offsets = - src_col.child(lists_column_view::offsets_column_index).data() + - src_col.offset(); - return d_offsets[src_list_id + 1] - d_offsets[src_list_id]; - }); - - // Compute offsets from sizes - thrust::exclusive_scan( - rmm::exec_policy(stream), d_out_offsets, d_out_offsets + num_lists + 1, d_out_offsets); - - return {std::move(list_offsets), std::move(validities)}; -} + data_type{type_id::INT32}, num_output_lists + 1, mask_state::UNALLOCATED, stream, mr); + auto const d_offsets = list_offsets->mutable_view().begin(); + + // The array of int8_t to store validities for list elements. + auto validities = rmm::device_uvector(has_null_mask ? num_output_lists : 0, stream); + + // Compute list sizes and validities. + thrust::transform( + rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(num_output_lists), + d_offsets, + [num_cols, + table_dv = *table_dv_ptr, + d_validities = validities.begin(), + has_null_mask] __device__(size_type const dst_list_id) { + auto const src_col_id = dst_list_id % num_cols; + auto const src_list_id = dst_list_id / num_cols; + auto const& src_lists_col = table_dv.column(src_col_id); + if (has_null_mask) { + d_validities[dst_list_id] = static_cast(src_lists_col.is_valid(src_list_id)); + } + auto const src_list_offsets = + src_lists_col.child(lists_column_view::offsets_column_index).data() + + src_lists_col.offset(); + return src_list_offsets[src_list_id + 1] - src_list_offsets[src_list_id]; + }); -/** - * @brief Creates child offsets, chars columns and null mask, null count of a strings column by - * applying the template function that can be used for computing the output size of each string as - * well as create the output. - * - * @tparam SizeAndExecuteFunction Function must accept an index and return a size. - * It must have members `d_offsets`, `d_chars`, and `d_validities` which are set to memory - * containing the offsets column, chars column and string validities during write. - * - * @param size_and_exec_fn This is called twice. Once for the output size of each string, which is - * written into the `d_offsets` array. After that, `d_chars` is set and this - * is called again to fill in the chars memory. The `d_validities` array may - * be modified to set the value `0` for the corresponding rows that contain - * null string elements. - * @param strings_count Number of strings. - * @param mr Device memory resource used to allocate the returned columns' device memory. - * @param stream CUDA stream used for device memory operations and kernel launches. - * @return offsets child column, chars child column, null_mask, and null_count for a strings column. - */ -template -std::tuple, std::unique_ptr, rmm::device_buffer, size_type> -make_strings_children_with_null_mask( - SizeAndExecuteFunction size_and_exec_fn, - size_type lists_count, - size_type strings_count, - - rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) -{ - auto offsets_column = make_numeric_column( - data_type{type_id::INT32}, strings_count + 1, mask_state::UNALLOCATED, stream, mr); - auto offsets_view = offsets_column->mutable_view(); - auto d_offsets = offsets_view.template data(); - size_and_exec_fn.d_offsets = d_offsets; - - auto validities = rmm::device_uvector(strings_count, stream); - size_and_exec_fn.d_validities = validities.begin(); - - // This is called twice: once for offsets and validities, and once for chars - auto for_each_fn = [lists_count, stream](SizeAndExecuteFunction& size_and_exec_fn) { - thrust::for_each_n(rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - lists_count, - size_and_exec_fn); - }; - - // Compute the string sizes (storing in `d_offsets`) and string validities - for_each_fn(size_and_exec_fn); - - printf("Line %d\n", __LINE__); - std::vector v(validities.size()); - CUDA_TRY(cudaMemcpyAsync( - v.data(), validities.data(), validities.size(), cudaMemcpyDefault, stream.value())); - stream.synchronize(); - for (auto x : v) { printf("%d, ", (int)x); } - printf("\n\n"); - - printf("Line %d\n", __LINE__); - cudf::test::print(offsets_view); - - // Compute the offsets from string sizes + // Compute offsets from sizes. thrust::exclusive_scan( - rmm::exec_policy(stream), d_offsets, d_offsets + strings_count + 1, d_offsets); - - printf("Line %d\n", __LINE__); - cudf::test::print(offsets_view); - - // Now build the chars column - auto const bytes = cudf::detail::get_value(offsets_view, strings_count, stream); - auto chars_column = - cudf::strings::detail::create_chars_child_column(strings_count, 0, bytes, stream, mr); - - // Execute the function fn again to fill the chars column. - // Note that if the output chars column has zero size, the function fn should not be called to - // avoid accidentally overwriting the offsets. - if (bytes > 0) { - size_and_exec_fn.d_chars = chars_column->mutable_view().template data(); - for_each_fn(size_and_exec_fn); - } - - // Finally compute null mask and null count from the validities array - auto [null_mask, null_count] = cudf::detail::valid_if( - validities.begin(), - validities.end(), - [] __device__(auto const valid) { return valid; }, - stream, - mr); + rmm::exec_policy(stream), d_offsets, d_offsets + num_output_lists + 1, d_offsets); - return std::make_tuple(std::move(offsets_column), - std::move(chars_column), - null_count > 0 ? std::move(null_mask) : rmm::device_buffer{}, - null_count); + return {std::move(list_offsets), std::move(validities)}; } /** - * @brief Compute string sizes, string validities, and concatenate strings functor. + * @brief Compute string sizes, string validities, and interleave string lists functor. * * This functor is executed twice. In the first pass, the sizes and validities of the output strings - * will be computed. In the second pass, this will concatenate the strings within each list element - * of the given lists column and apply the separator. The null-replacement string scalar - * `string_narep_dv` (if valid) is used in place of any null string. - * - * @tparam Functor The functor which can check for validity of the input list at a given list index - * as well as access to the separator corresponding to the list index. + * will be computed. In the second pass, this will interleave the lists of strings of the given + * table of lists columns. */ -struct compute_size_and_copy_fn { +struct compute_size_and_interleave_lists_fn { table_device_view const table_dv; - // Store offsets of the lists in the output lists column + // Store list offsets of the output lists column. offset_type const* const dst_list_offsets; + // Flag to specify whether to compute string validities. bool const has_null_mask; - // Store offsets of the strings in the chars column + // Store offsets of the strings in the chars column. offset_type* d_offsets{nullptr}; // If d_chars == nullptr: only compute sizes and validities of the output strings. @@ -225,10 +120,7 @@ struct compute_size_and_copy_fn { auto const src_list_id = dst_list_id / num_cols; auto const& src_lists_col = table_dv.column(src_col_id); - if (has_null_mask and src_lists_col.is_null(src_list_id)) { - printf("return c null\n"); - return; - } + if (has_null_mask and src_lists_col.is_null(src_list_id)) { return; } auto const src_list_offsets = src_lists_col.child(lists_column_view::offsets_column_index).data() + @@ -244,11 +136,12 @@ struct compute_size_and_copy_fn { end_idx = src_list_offsets[src_list_id + 1]; read_idx < end_idx; ++read_idx, ++write_idx) { - auto const is_valid = src_child.is_valid(read_idx); - if (has_null_mask) { d_validities[write_idx] = static_cast(is_valid); } + if (has_null_mask) { + d_validities[write_idx] = static_cast(src_child.is_valid(read_idx)); + } d_offsets[write_idx] = src_child_offsets[read_idx + 1] - src_child_offsets[read_idx]; } - } else { // just copy the entire list of strings + } else { // just copy the entire memory region containing all strings in the list auto const start_idx = src_child_offsets[src_list_offsets[src_list_id]]; auto const end_idx = src_child_offsets[src_list_offsets[src_list_id + 1]]; if (start_idx < end_idx) { @@ -263,13 +156,13 @@ struct compute_size_and_copy_fn { }; /** - * @brief Struct used in type_dispatcher to copy entries to the output lists column + * @brief Struct used in type_dispatcher to copy entries to the output lists column. */ struct copy_list_entries_fn { template std::enable_if_t, std::unique_ptr> operator()( table_view const& input, - column_view const& list_offsets, + column_view const& output_list_offsets, size_type num_output_lists, size_type num_output_entries, bool has_null_mask, @@ -277,35 +170,16 @@ struct copy_list_entries_fn { rmm::mr::device_memory_resource* mr) const noexcept { auto const table_dv_ptr = table_device_view::create(input); - auto const comp_fn = - compute_size_and_copy_fn{*table_dv_ptr, list_offsets.begin(), has_null_mask}; - // if (not has_null_mask) { - // auto [offsets_column, chars_column] = - // cudf::strings::detail::make_strings_children(comp_fn, num_output_lists, 0, stream, - // mr); - // return make_strings_column(num_output_entries, - // std::move(offsets_column), - // std::move(chars_column), - // 0, - // rmm::device_buffer{}, - // stream, - // mr); - // } - - printf("numlist: %d, entri: %d\n", num_output_lists, num_output_entries); + auto const comp_fn = compute_size_and_interleave_lists_fn{ + *table_dv_ptr, output_list_offsets.begin(), has_null_mask}; + // Cannot call `cudf::strings::detail::make_strings_children` because that function executes the + // functor `comp_fn` on the range equal to `num_output_entries`. auto [offsets_column, chars_column, null_mask, null_count] = - make_strings_children_with_null_mask( + cudf::strings::detail::make_strings_children_with_null_mask( comp_fn, num_output_lists, num_output_entries, stream, mr); - if (not has_null_mask) - return make_strings_column(num_output_entries, - std::move(offsets_column), - std::move(chars_column), - 0, - rmm::device_buffer{}, - stream, - mr); - else + + if (has_null_mask) { return make_strings_column(num_output_entries, std::move(offsets_column), std::move(chars_column), @@ -313,30 +187,38 @@ struct copy_list_entries_fn { std::move(null_mask), stream, mr); + } + return make_strings_column(num_output_entries, + std::move(offsets_column), + std::move(chars_column), + 0, + rmm::device_buffer{}, + stream, + mr); } template std::enable_if_t(), std::unique_ptr> operator()( table_view const& input, - column_view const& list_offsets, + column_view const& output_list_offsets, size_type num_output_lists, size_type num_output_entries, bool has_null_mask, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) const noexcept { - auto const child_col = lists_column_view(*input.begin()).child(); auto const num_cols = input.num_columns(); auto const num_rows = input.num_rows(); auto const table_dv_ptr = table_device_view::create(input); + // The output child column. + auto const child_col = lists_column_view(*input.begin()).child(); auto output = allocate_like(child_col, num_output_entries, mask_allocation_policy::NEVER, stream, mr); auto output_dv_ptr = mutable_column_device_view::create(*output); - // The array of int8_t to store element validities - auto validities = has_null_mask ? rmm::device_uvector(num_output_entries, stream) - : rmm::device_uvector(0, stream); + // The array of int8_t to store entry validities. + auto validities = rmm::device_uvector(has_null_mask ? num_output_entries : 0, stream); thrust::for_each_n( rmm::exec_policy(stream), @@ -344,9 +226,9 @@ struct copy_list_entries_fn { num_output_lists, [num_cols, table_dv = *table_dv_ptr, - out_validities = validities.begin(), - dst_list_offsets = list_offsets.begin(), - out_entries = output_dv_ptr->begin(), + d_validities = validities.begin(), + dst_list_offsets = output_list_offsets.begin(), + d_output = output_dv_ptr->begin(), has_null_mask] __device__(size_type const dst_list_id) { auto const src_col_id = dst_list_id % num_cols; auto const src_list_id = dst_list_id / num_cols; @@ -355,15 +237,27 @@ struct copy_list_entries_fn { src_lists_col.child(lists_column_view::offsets_column_index).data() + src_lists_col.offset(); auto const& src_child = src_lists_col.child(lists_column_view::child_column_index); - for (auto read_idx = src_list_offsets[src_list_id], - end_idx = src_list_offsets[src_list_id + 1], - write_idx = dst_list_offsets[dst_list_id]; - read_idx < end_idx; - ++read_idx, ++write_idx) { - auto const is_valid = src_child.is_valid(read_idx); - if (has_null_mask) { out_validities[write_idx] = static_cast(is_valid); } - out_entries[write_idx] = is_valid ? src_child.element(read_idx) : T{}; + + // The indices of the entries in the list to gather from. + auto const start_idx = src_list_offsets[src_list_id]; + auto const end_idx = src_list_offsets[src_list_id + 1]; + + // Fill the validities array if necessary. + if (has_null_mask) { + for (auto read_idx = src_list_offsets[src_list_id], + end_idx = src_list_offsets[src_list_id + 1], + write_idx = dst_list_offsets[dst_list_id]; + read_idx < end_idx; + ++read_idx, ++write_idx) { + auto const is_valid = src_child.is_valid(read_idx); + d_validities[write_idx] = static_cast(is_valid); + } } + + // Do a memcpy for the entire list entries. + auto const input_ptr = src_child.data() + start_idx; + auto const output_ptr = &d_output[dst_list_offsets[dst_list_id]]; + std::memcpy(output_ptr, input_ptr, sizeof(T) * (end_idx - start_idx)); }); if (has_null_mask) { @@ -417,22 +311,23 @@ std::unique_ptr interleave_columns(table_view const& input, "The types of entries in the input columns must be the same."); } - if (input.num_columns() == 1) { return std::make_unique(*(input.begin()), stream, mr); } if (input.num_rows() == 0) { return cudf::empty_like(input.column(0)); } + if (input.num_columns() == 1) { return std::make_unique(*(input.begin()), stream, mr); } // Generate offsets of the output lists column auto [list_offsets, list_validities] = generate_list_offsets_and_validities(input, has_null_mask, stream, mr); + auto const offsets_view = list_offsets->view(); // Copy entries from the input lists columns to the output lists column - this needed to be // specialized for different types auto const num_output_lists = input.num_rows() * input.num_columns(); auto const num_output_entries = - cudf::detail::get_value(list_offsets->view(), num_output_lists, stream); + cudf::detail::get_value(offsets_view, num_output_lists, stream); auto list_entries = type_dispatcher(entry_type, copy_list_entries_fn{}, input, - list_offsets->view(), + offsets_view, num_output_lists, num_output_entries, has_null_mask, From 48a6b147e85b9e94143437f9c4352984d5e7c1db Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Fri, 23 Apr 2021 11:04:07 -0600 Subject: [PATCH 26/36] Use `thrust::identity` and `thrust::copy` --- cpp/src/lists/interleave_columns.cu | 22 ++++++++-------------- 1 file changed, 8 insertions(+), 14 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 47cb6be4acb..18c1ae1a4d5 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -27,6 +27,7 @@ #include #include +#include #include namespace cudf { @@ -149,7 +150,7 @@ struct compute_string_sizes_and_interleave_lists_fn { src_child.child(strings_column_view::chars_column_index).data() + src_child.offset() + start_idx; auto const output_ptr = d_chars + d_offsets[write_idx]; - std::memcpy(output_ptr, input_ptr, end_idx - start_idx); + thrust::copy(thrust::seq, input_ptr, input_ptr + end_idx - start_idx, output_ptr); } } } @@ -256,18 +257,15 @@ struct interleave_list_entries_fn { } // Do a memcpy for the entire list entries. - auto const input_ptr = src_child.data() + start_idx; - auto const output_ptr = &d_output[dst_list_offsets[dst_list_id]]; - std::memcpy(output_ptr, input_ptr, sizeof(T) * (end_idx - start_idx)); + auto const input_ptr = reinterpret_cast(src_child.data() + start_idx); + auto const output_ptr = reinterpret_cast(&d_output[dst_list_offsets[dst_list_id]]); + thrust::copy( + thrust::seq, input_ptr, input_ptr + sizeof(T) * (end_idx - start_idx), output_ptr); }); if (has_null_mask) { auto [null_mask, null_count] = cudf::detail::valid_if( - validities.begin(), - validities.end(), - [] __device__(auto const valid) { return valid; }, - stream, - mr); + validities.begin(), validities.end(), thrust::identity{}, stream, mr); if (null_count > 0) { output->set_null_mask(null_mask, null_count); } } @@ -346,11 +344,7 @@ std::unique_ptr interleave_columns(table_view const& input, } auto [null_mask, null_count] = cudf::detail::valid_if( - list_validities.begin(), - list_validities.end(), - [] __device__(auto const valid) { return valid; }, - stream, - mr); + list_validities.begin(), list_validities.end(), thrust::identity{}, stream, mr); return make_lists_column(num_output_lists, std::move(list_offsets), std::move(list_entries), From 9d66ce228c672f65682cf3f35eb57c74a2c7d4f9 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Tue, 27 Apr 2021 20:36:11 -0600 Subject: [PATCH 27/36] Edit comments --- cpp/src/lists/interleave_columns.cu | 25 ++++++++++++------------- 1 file changed, 12 insertions(+), 13 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 18c1ae1a4d5..cc7a20d2946 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -131,6 +131,7 @@ struct compute_string_sizes_and_interleave_lists_fn { src_child.child(strings_column_view::offsets_column_index).data() + src_child.offset(); + // read_idx and write_idx are indices of string elements. size_type write_idx = dst_list_offsets[dst_list_id]; if (not d_chars) { // just compute sizes of strings within a list for (auto read_idx = src_list_offsets[src_list_id], @@ -143,6 +144,7 @@ struct compute_string_sizes_and_interleave_lists_fn { d_offsets[write_idx] = src_child_offsets[read_idx + 1] - src_child_offsets[read_idx]; } } else { // just copy the entire memory region containing all strings in the list + // start_idx and end_idx are indices of character elements. auto const start_idx = src_child_offsets[src_list_offsets[src_list_id]]; auto const end_idx = src_child_offsets[src_list_offsets[src_list_id + 1]]; if (start_idx < end_idx) { @@ -181,7 +183,7 @@ struct interleave_list_entries_fn { cudf::strings::detail::make_strings_children_with_null_mask( comp_fn, num_output_lists, num_output_entries, stream, mr); - if (has_null_mask) { + if (has_null_mask and null_count > 0) { return make_strings_column(num_output_entries, std::move(offsets_column), std::move(chars_column), @@ -241,24 +243,21 @@ struct interleave_list_entries_fn { auto const& src_child = src_lists_col.child(lists_column_view::child_column_index); // The indices of the entries in the list to gather from. - auto const start_idx = src_list_offsets[src_list_id]; - auto const end_idx = src_list_offsets[src_list_id + 1]; + auto const start_idx = src_list_offsets[src_list_id]; + auto const end_idx = src_list_offsets[src_list_id + 1]; + auto const write_start = dst_list_offsets[dst_list_id]; // Fill the validities array if necessary. if (has_null_mask) { - for (auto read_idx = src_list_offsets[src_list_id], - end_idx = src_list_offsets[src_list_id + 1], - write_idx = dst_list_offsets[dst_list_id]; - read_idx < end_idx; + for (auto read_idx = start_idx, write_idx = write_start; read_idx < end_idx; ++read_idx, ++write_idx) { - auto const is_valid = src_child.is_valid(read_idx); - d_validities[write_idx] = static_cast(is_valid); + d_validities[write_idx] = static_cast(src_child.is_valid(read_idx)); } } - // Do a memcpy for the entire list entries. + // Do a copy for the entire list entries. auto const input_ptr = reinterpret_cast(src_child.data() + start_idx); - auto const output_ptr = reinterpret_cast(&d_output[dst_list_offsets[dst_list_id]]); + auto const output_ptr = reinterpret_cast(&d_output[write_start]); thrust::copy( thrust::seq, input_ptr, input_ptr + sizeof(T) * (end_idx - start_idx), output_ptr); }); @@ -313,13 +312,13 @@ std::unique_ptr interleave_columns(table_view const& input, if (input.num_rows() == 0) { return cudf::empty_like(input.column(0)); } if (input.num_columns() == 1) { return std::make_unique(*(input.begin()), stream, mr); } - // Generate offsets of the output lists column + // Generate offsets of the output lists column. auto [list_offsets, list_validities] = generate_list_offsets_and_validities(input, has_null_mask, stream, mr); auto const offsets_view = list_offsets->view(); // Copy entries from the input lists columns to the output lists column - this needed to be - // specialized for different types + // specialized for different types. auto const num_output_lists = input.num_rows() * input.num_columns(); auto const num_output_entries = cudf::detail::get_value(offsets_view, num_output_lists, stream); From 467845f12124f2f582feb764aa22900a1aa1dc98 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 28 Apr 2021 07:55:28 -0600 Subject: [PATCH 28/36] Fix doxygen --- cpp/include/cudf/lists/detail/interleave_columns.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/cudf/lists/detail/interleave_columns.hpp b/cpp/include/cudf/lists/detail/interleave_columns.hpp index 831ca5d885b..7ae90779fdc 100644 --- a/cpp/include/cudf/lists/detail/interleave_columns.hpp +++ b/cpp/include/cudf/lists/detail/interleave_columns.hpp @@ -39,7 +39,7 @@ namespace detail { * @throws cudf::logic_error if all lists columns do not have the same entry type. * * @param input Table containing lists columns to interleave. - * @param has_null_mask A boolean flag indicating that the input columns have null mask. + * @param has_null_mask A boolean flag indicating that the input columns have a null mask. * @param stream CUDA stream used for device memory operations and kernel launches. * @param mr Device memory resource used to allocate the returned column's device memory. * @return The interleaved columns as a single column. From 53dda6d796870fc974b377c965b44fec4b83fa17 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 28 Apr 2021 07:55:44 -0600 Subject: [PATCH 29/36] Change `size_type` to `offset_type` --- cpp/src/lists/interleave_columns.cu | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index cc7a20d2946..8ae16f6fdcf 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -76,7 +76,7 @@ generate_list_offsets_and_validities(table_view const& input, d_validities[dst_list_id] = static_cast(src_lists_col.is_valid(src_list_id)); } auto const src_list_offsets = - src_lists_col.child(lists_column_view::offsets_column_index).data() + + src_lists_col.child(lists_column_view::offsets_column_index).data() + src_lists_col.offset(); return src_list_offsets[src_list_id + 1] - src_list_offsets[src_list_id]; }); @@ -124,11 +124,11 @@ struct compute_string_sizes_and_interleave_lists_fn { if (has_null_mask and src_lists_col.is_null(src_list_id)) { return; } auto const src_list_offsets = - src_lists_col.child(lists_column_view::offsets_column_index).data() + + src_lists_col.child(lists_column_view::offsets_column_index).data() + src_lists_col.offset(); auto const& src_child = src_lists_col.child(lists_column_view::child_column_index); auto const src_child_offsets = - src_child.child(strings_column_view::offsets_column_index).data() + + src_child.child(strings_column_view::offsets_column_index).data() + src_child.offset(); // read_idx and write_idx are indices of string elements. @@ -238,7 +238,7 @@ struct interleave_list_entries_fn { auto const src_list_id = dst_list_id / num_cols; auto const& src_lists_col = table_dv.column(src_col_id); auto const src_list_offsets = - src_lists_col.child(lists_column_view::offsets_column_index).data() + + src_lists_col.child(lists_column_view::offsets_column_index).data() + src_lists_col.offset(); auto const& src_child = src_lists_col.child(lists_column_view::child_column_index); @@ -321,7 +321,7 @@ std::unique_ptr interleave_columns(table_view const& input, // specialized for different types. auto const num_output_lists = input.num_rows() * input.num_columns(); auto const num_output_entries = - cudf::detail::get_value(offsets_view, num_output_lists, stream); + cudf::detail::get_value(offsets_view, num_output_lists, stream); auto list_entries = type_dispatcher(entry_type, interleave_list_entries_fn{}, input, From 0ccb82f85c6a2125e5f5e7c21160e1c97c449a50 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 28 Apr 2021 10:24:56 -0600 Subject: [PATCH 30/36] Rename variables and some other improvements --- cpp/src/lists/interleave_columns.cu | 125 ++++++++++++++-------------- 1 file changed, 61 insertions(+), 64 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 8ae16f6fdcf..047af54db7f 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -68,17 +68,15 @@ generate_list_offsets_and_validities(table_view const& input, [num_cols, table_dv = *table_dv_ptr, d_validities = validities.begin(), - has_null_mask] __device__(size_type const dst_list_id) { - auto const src_col_id = dst_list_id % num_cols; - auto const src_list_id = dst_list_id / num_cols; - auto const& src_lists_col = table_dv.column(src_col_id); - if (has_null_mask) { - d_validities[dst_list_id] = static_cast(src_lists_col.is_valid(src_list_id)); - } - auto const src_list_offsets = - src_lists_col.child(lists_column_view::offsets_column_index).data() + - src_lists_col.offset(); - return src_list_offsets[src_list_id + 1] - src_list_offsets[src_list_id]; + has_null_mask] __device__(size_type const idx) { + auto const col_id = idx % num_cols; + auto const list_id = idx / num_cols; + auto const& lists_col = table_dv.column(col_id); + if (has_null_mask) { d_validities[idx] = static_cast(lists_col.is_valid(list_id)); } + auto const list_offsets = + lists_col.child(lists_column_view::offsets_column_index).data() + + lists_col.offset(); + return list_offsets[list_id + 1] - list_offsets[list_id]; }); // Compute offsets from sizes. @@ -93,7 +91,7 @@ generate_list_offsets_and_validities(table_view const& input, * * This functor is executed twice. In the first pass, the sizes and validities of the output strings * will be computed. In the second pass, this will interleave the lists of strings of the given - * table of lists columns. + * table containing those lists. */ struct compute_string_sizes_and_interleave_lists_fn { table_device_view const table_dv; @@ -104,7 +102,7 @@ struct compute_string_sizes_and_interleave_lists_fn { // Flag to specify whether to compute string validities. bool const has_null_mask; - // Store offsets of the strings in the chars column. + // Store offsets of the strings. offset_type* d_offsets{nullptr}; // If d_chars == nullptr: only compute sizes and validities of the output strings. @@ -114,43 +112,40 @@ struct compute_string_sizes_and_interleave_lists_fn { // We need to set `1` or `0` for the validities of the strings in the child column. int8_t* d_validities{nullptr}; - __device__ void operator()(size_type const dst_list_id) + __device__ void operator()(size_type const idx) { - auto const num_cols = table_dv.num_columns(); - auto const src_col_id = dst_list_id % num_cols; - auto const src_list_id = dst_list_id / num_cols; + auto const num_cols = table_dv.num_columns(); + auto const col_id = idx % num_cols; + auto const list_id = idx / num_cols; - auto const& src_lists_col = table_dv.column(src_col_id); - if (has_null_mask and src_lists_col.is_null(src_list_id)) { return; } + auto const& lists_col = table_dv.column(col_id); + if (has_null_mask and lists_col.is_null(list_id)) { return; } - auto const src_list_offsets = - src_lists_col.child(lists_column_view::offsets_column_index).data() + - src_lists_col.offset(); - auto const& src_child = src_lists_col.child(lists_column_view::child_column_index); - auto const src_child_offsets = - src_child.child(strings_column_view::offsets_column_index).data() + - src_child.offset(); + auto const list_offsets = + lists_col.child(lists_column_view::offsets_column_index).data() + + lists_col.offset(); + auto const& str_col = lists_col.child(lists_column_view::child_column_index); + auto const str_offsets = + str_col.child(strings_column_view::offsets_column_index).data(); // read_idx and write_idx are indices of string elements. - size_type write_idx = dst_list_offsets[dst_list_id]; + size_type write_idx = dst_list_offsets[idx]; if (not d_chars) { // just compute sizes of strings within a list - for (auto read_idx = src_list_offsets[src_list_id], - end_idx = src_list_offsets[src_list_id + 1]; + for (auto read_idx = list_offsets[list_id], end_idx = list_offsets[list_id + 1]; read_idx < end_idx; ++read_idx, ++write_idx) { if (has_null_mask) { - d_validities[write_idx] = static_cast(src_child.is_valid(read_idx)); + d_validities[write_idx] = static_cast(str_col.is_valid(read_idx)); } - d_offsets[write_idx] = src_child_offsets[read_idx + 1] - src_child_offsets[read_idx]; + d_offsets[write_idx] = str_offsets[read_idx + 1] - str_offsets[read_idx]; } } else { // just copy the entire memory region containing all strings in the list // start_idx and end_idx are indices of character elements. - auto const start_idx = src_child_offsets[src_list_offsets[src_list_id]]; - auto const end_idx = src_child_offsets[src_list_offsets[src_list_id + 1]]; + auto const start_idx = str_offsets[list_offsets[list_id]]; + auto const end_idx = str_offsets[list_offsets[list_id + 1]]; if (start_idx < end_idx) { auto const input_ptr = - src_child.child(strings_column_view::chars_column_index).data() + - src_child.offset() + start_idx; + str_col.child(strings_column_view::chars_column_index).data() + start_idx; auto const output_ptr = d_chars + d_offsets[write_idx]; thrust::copy(thrust::seq, input_ptr, input_ptr + end_idx - start_idx, output_ptr); } @@ -177,13 +172,10 @@ struct interleave_list_entries_fn { auto const comp_fn = compute_string_sizes_and_interleave_lists_fn{ *table_dv_ptr, output_list_offsets.begin(), has_null_mask}; - // Cannot call `cudf::strings::detail::make_strings_children` because that function executes the - // functor `comp_fn` on the range equal to `num_output_entries`. - auto [offsets_column, chars_column, null_mask, null_count] = - cudf::strings::detail::make_strings_children_with_null_mask( - comp_fn, num_output_lists, num_output_entries, stream, mr); - - if (has_null_mask and null_count > 0) { + if (has_null_mask) { + auto [offsets_column, chars_column, null_mask, null_count] = + cudf::strings::detail::make_strings_children_with_null_mask( + comp_fn, num_output_lists, num_output_entries, stream, mr); return make_strings_column(num_output_entries, std::move(offsets_column), std::move(chars_column), @@ -192,6 +184,9 @@ struct interleave_list_entries_fn { stream, mr); } + + auto [offsets_column, chars_column] = cudf::strings::detail::make_strings_children( + comp_fn, num_output_lists, num_output_entries, stream, mr); return make_strings_column(num_output_entries, std::move(offsets_column), std::move(chars_column), @@ -216,9 +211,11 @@ struct interleave_list_entries_fn { auto const table_dv_ptr = table_device_view::create(input); // The output child column. - auto const child_col = lists_column_view(*input.begin()).child(); - auto output = - allocate_like(child_col, num_output_entries, mask_allocation_policy::NEVER, stream, mr); + auto output = allocate_like(lists_column_view(*input.begin()).child(), + num_output_entries, + mask_allocation_policy::NEVER, + stream, + mr); auto output_dv_ptr = mutable_column_device_view::create(*output); // The array of int8_t to store entry validities. @@ -229,34 +226,34 @@ struct interleave_list_entries_fn { thrust::make_counting_iterator(0), num_output_lists, [num_cols, - table_dv = *table_dv_ptr, - d_validities = validities.begin(), - dst_list_offsets = output_list_offsets.begin(), - d_output = output_dv_ptr->begin(), - has_null_mask] __device__(size_type const dst_list_id) { - auto const src_col_id = dst_list_id % num_cols; - auto const src_list_id = dst_list_id / num_cols; - auto const& src_lists_col = table_dv.column(src_col_id); - auto const src_list_offsets = - src_lists_col.child(lists_column_view::offsets_column_index).data() + - src_lists_col.offset(); - auto const& src_child = src_lists_col.child(lists_column_view::child_column_index); - - // The indices of the entries in the list to gather from. - auto const start_idx = src_list_offsets[src_list_id]; - auto const end_idx = src_list_offsets[src_list_id + 1]; - auto const write_start = dst_list_offsets[dst_list_id]; + table_dv = *table_dv_ptr, + d_validities = validities.begin(), + d_offsets = output_list_offsets.begin(), + d_output = output_dv_ptr->begin(), + has_null_mask] __device__(size_type const idx) { + auto const col_id = idx % num_cols; + auto const list_id = idx / num_cols; + auto const& lists_col = table_dv.column(col_id); + auto const list_offsets = + lists_col.child(lists_column_view::offsets_column_index).data() + + lists_col.offset(); + auto const& data_col = lists_col.child(lists_column_view::child_column_index); + + // The indices of the entries within the source list. + auto const start_idx = list_offsets[list_id]; + auto const end_idx = list_offsets[list_id + 1]; + auto const write_start = d_offsets[idx]; // Fill the validities array if necessary. if (has_null_mask) { for (auto read_idx = start_idx, write_idx = write_start; read_idx < end_idx; ++read_idx, ++write_idx) { - d_validities[write_idx] = static_cast(src_child.is_valid(read_idx)); + d_validities[write_idx] = static_cast(data_col.is_valid(read_idx)); } } // Do a copy for the entire list entries. - auto const input_ptr = reinterpret_cast(src_child.data() + start_idx); + auto const input_ptr = reinterpret_cast(data_col.data() + start_idx); auto const output_ptr = reinterpret_cast(&d_output[write_start]); thrust::copy( thrust::seq, input_ptr, input_ptr + sizeof(T) * (end_idx - start_idx), output_ptr); From ed45140a0680a3fc6ee42d7c399605fc7c85c4fa Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 28 Apr 2021 10:40:01 -0600 Subject: [PATCH 31/36] Rewrite unit tests --- .../reshape/interleave_columns_tests.cpp | 83 +++++++++++++------ 1 file changed, 59 insertions(+), 24 deletions(-) diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index 2f21ebaa082..5783eeb4526 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -426,8 +426,9 @@ struct ListsColumnsInterleaveTypedTest : public cudf::test::BaseFixture { }; #define ListsCol cudf::test::lists_column_wrapper -using TypesForTest = - cudf::test::Concat; +using TypesForTest = cudf::test::Concat; TYPED_TEST_CASE(ListsColumnsInterleaveTypedTest, TypesForTest); TYPED_TEST(ListsColumnsInterleaveTypedTest, InterleaveEmptyColumns) @@ -511,7 +512,6 @@ TYPED_TEST(ListsColumnsInterleaveTypedTest, SimpleInputWithNulls) ListsCol{{null, null, null, null, null}, all_nulls()}}, null_at(0)} .release(); - auto const expected = ListsCol{{ListsCol{{1, null, 3, 4}, null_at(1)}, ListsCol{{10, 11, 12, null}, null_at(3)}, ListsCol{} /*NULL*/, @@ -567,12 +567,12 @@ TEST_F(ListsColumnsInterleaveTest, SimpleInputStringsColumnsWithNulls) TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputNoNull) { - auto const col_original = ListsCol{{1, 2, 3}, {2, 3}, {3, 4, 5, 6}, {5, 6}, {}, {7}}.release(); - auto const col1 = cudf::slice(col_original->view(), {0, 3})[0]; - auto const col2 = cudf::slice(col_original->view(), {1, 4})[0]; - auto const col3 = cudf::slice(col_original->view(), {2, 5})[0]; - auto const col4 = cudf::slice(col_original->view(), {3, 6})[0]; - auto const expected = ListsCol{ + auto const col = ListsCol{{1, 2, 3}, {2, 3}, {3, 4, 5, 6}, {5, 6}, {}, {7}}.release(); + auto const col1 = cudf::slice(col->view(), {0, 3})[0]; + auto const col2 = cudf::slice(col->view(), {1, 4})[0]; + auto const col3 = cudf::slice(col->view(), {2, 5})[0]; + auto const col4 = cudf::slice(col->view(), {3, 6})[0]; + auto const expected = ListsCol{ ListsCol{1, 2, 3}, ListsCol{2, 3}, ListsCol{3, 4, 5, 6}, @@ -591,21 +591,20 @@ TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputNoNull) TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputWithNulls) { - auto const col_original = ListsCol{{ListsCol{{null, 2, 3}, null_at(0)}, - ListsCol{2, 3}, /*NULL*/ - ListsCol{{3, null, 5, 6}, null_at(1)}, - ListsCol{5, 6}, /*NULL*/ - ListsCol{}, /*NULL*/ - ListsCol{7}, - ListsCol{8, 9, 10}}, - null_at({1, 3, 4})} - .release(); - - auto const col1 = cudf::slice(col_original->view(), {0, 3})[0]; - auto const col2 = cudf::slice(col_original->view(), {1, 4})[0]; - auto const col3 = cudf::slice(col_original->view(), {2, 5})[0]; - auto const col4 = cudf::slice(col_original->view(), {3, 6})[0]; - auto const col5 = cudf::slice(col_original->view(), {4, 7})[0]; + auto const col = ListsCol{{ListsCol{{null, 2, 3}, null_at(0)}, + ListsCol{2, 3}, /*NULL*/ + ListsCol{{3, null, 5, 6}, null_at(1)}, + ListsCol{5, 6}, /*NULL*/ + ListsCol{}, /*NULL*/ + ListsCol{7}, + ListsCol{8, 9, 10}}, + null_at({1, 3, 4})} + .release(); + auto const col1 = cudf::slice(col->view(), {0, 3})[0]; + auto const col2 = cudf::slice(col->view(), {1, 4})[0]; + auto const col3 = cudf::slice(col->view(), {2, 5})[0]; + auto const col4 = cudf::slice(col->view(), {3, 6})[0]; + auto const col5 = cudf::slice(col->view(), {4, 7})[0]; auto const expected = ListsCol{{ListsCol{{null, 2, 3}, null_at(0)}, ListsCol{}, /*NULL*/ ListsCol{{3, null, 5, 6}, null_at(1)}, @@ -627,4 +626,40 @@ TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputWithNulls) CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); } +TEST_F(ListsColumnsInterleaveTest, SlicedStringsColumnsInputWithNulls) +{ + auto const col = + StrListsCol{ + {StrListsCol{{"Tomato", "Bear" /*NULL*/, "Apple"}, null_at(1)}, + StrListsCol{{"Banana", "Pig" /*NULL*/, "Kiwi", "Cherry", "Whale" /*NULL*/}, null_at({1, 4})}, + StrListsCol{"Coconut"}, + StrListsCol{{"Orange", "Dog" /*NULL*/, "Fox" /*NULL*/, "Duck" /*NULL*/}, null_at({1, 2, 3})}, + StrListsCol{"Lemon", "Peach"}, + StrListsCol{{"Deer" /*NULL*/, "Snake" /*NULL*/, "Horse" /*NULL*/}, all_nulls()}}, /*NULL*/ + null_at(5)} + .release(); + auto const col1 = cudf::slice(col->view(), {0, 3})[0]; + auto const col2 = cudf::slice(col->view(), {1, 4})[0]; + auto const col3 = cudf::slice(col->view(), {2, 5})[0]; + auto const col4 = cudf::slice(col->view(), {3, 6})[0]; + auto const expected = + StrListsCol{ + {StrListsCol{{"Tomato", "" /*NULL*/, "Apple"}, null_at(1)}, + StrListsCol{{"Banana", "" /*NULL*/, "Kiwi", "Cherry", "" /*NULL*/}, null_at({1, 4})}, + StrListsCol{"Coconut"}, + StrListsCol{{"Orange", "" /*NULL*/, "" /*NULL*/, "" /*NULL*/}, null_at({1, 2, 3})}, + StrListsCol{{"Banana", "" /*NULL*/, "Kiwi", "Cherry", "" /*NULL*/}, null_at({1, 4})}, + StrListsCol{"Coconut"}, + StrListsCol{{"Orange", "" /*NULL*/, "" /*NULL*/, "" /*NULL*/}, null_at({1, 2, 3})}, + StrListsCol{"Lemon", "Peach"}, + StrListsCol{"Coconut"}, + StrListsCol{{"Orange", "" /*NULL*/, "" /*NULL*/, "" /*NULL*/}, null_at({1, 2, 3})}, + StrListsCol{"Lemon", "Peach"}, + StrListsCol{}}, /*NULL*/ + null_at(11)} + .release(); + auto const results = cudf::interleave_columns(TView{{col1, col2, col3, col4}}); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*expected, *results, print_all); +} + CUDF_TEST_PROGRAM_MAIN() From d22c3170d9bb102cf9260154f5752721eb7054fa Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 28 Apr 2021 11:06:44 -0600 Subject: [PATCH 32/36] Add keyword `template` to template function calls --- cpp/src/lists/interleave_columns.cu | 17 +++++++++-------- 1 file changed, 9 insertions(+), 8 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 047af54db7f..106efd73de3 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -74,7 +74,7 @@ generate_list_offsets_and_validities(table_view const& input, auto const& lists_col = table_dv.column(col_id); if (has_null_mask) { d_validities[idx] = static_cast(lists_col.is_valid(list_id)); } auto const list_offsets = - lists_col.child(lists_column_view::offsets_column_index).data() + + lists_col.child(lists_column_view::offsets_column_index).template data() + lists_col.offset(); return list_offsets[list_id + 1] - list_offsets[list_id]; }); @@ -122,11 +122,11 @@ struct compute_string_sizes_and_interleave_lists_fn { if (has_null_mask and lists_col.is_null(list_id)) { return; } auto const list_offsets = - lists_col.child(lists_column_view::offsets_column_index).data() + + lists_col.child(lists_column_view::offsets_column_index).template data() + lists_col.offset(); auto const& str_col = lists_col.child(lists_column_view::child_column_index); auto const str_offsets = - str_col.child(strings_column_view::offsets_column_index).data(); + str_col.child(strings_column_view::offsets_column_index).template data(); // read_idx and write_idx are indices of string elements. size_type write_idx = dst_list_offsets[idx]; @@ -145,7 +145,7 @@ struct compute_string_sizes_and_interleave_lists_fn { auto const end_idx = str_offsets[list_offsets[list_id + 1]]; if (start_idx < end_idx) { auto const input_ptr = - str_col.child(strings_column_view::chars_column_index).data() + start_idx; + str_col.child(strings_column_view::chars_column_index).template data() + start_idx; auto const output_ptr = d_chars + d_offsets[write_idx]; thrust::copy(thrust::seq, input_ptr, input_ptr + end_idx - start_idx, output_ptr); } @@ -228,14 +228,14 @@ struct interleave_list_entries_fn { [num_cols, table_dv = *table_dv_ptr, d_validities = validities.begin(), - d_offsets = output_list_offsets.begin(), - d_output = output_dv_ptr->begin(), + d_offsets = output_list_offsets.template begin(), + d_output = output_dv_ptr->template begin(), has_null_mask] __device__(size_type const idx) { auto const col_id = idx % num_cols; auto const list_id = idx / num_cols; auto const& lists_col = table_dv.column(col_id); auto const list_offsets = - lists_col.child(lists_column_view::offsets_column_index).data() + + lists_col.child(lists_column_view::offsets_column_index).template data() + lists_col.offset(); auto const& data_col = lists_col.child(lists_column_view::child_column_index); @@ -253,7 +253,8 @@ struct interleave_list_entries_fn { } // Do a copy for the entire list entries. - auto const input_ptr = reinterpret_cast(data_col.data() + start_idx); + auto const input_ptr = + reinterpret_cast(data_col.template data() + start_idx); auto const output_ptr = reinterpret_cast(&d_output[write_start]); thrust::copy( thrust::seq, input_ptr, input_ptr + sizeof(T) * (end_idx - start_idx), output_ptr); From 0e50dfce02fd08a17ca017828e1b223cfe4807dc Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 28 Apr 2021 11:08:16 -0600 Subject: [PATCH 33/36] Add header to conda recipes --- conda/recipes/libcudf/meta.yaml | 1 + 1 file changed, 1 insertion(+) diff --git a/conda/recipes/libcudf/meta.yaml b/conda/recipes/libcudf/meta.yaml index 75955428eab..7f02aee819a 100644 --- a/conda/recipes/libcudf/meta.yaml +++ b/conda/recipes/libcudf/meta.yaml @@ -135,6 +135,7 @@ test: - test -f $PREFIX/include/cudf/lists/detail/concatenate.hpp - test -f $PREFIX/include/cudf/lists/detail/copying.hpp - test -f $PREFIX/include/cudf/lists/detail/drop_list_duplicates.hpp + - test -f $PREFIX/include/cudf/lists/detail/interleave_columns.hpp - test -f $PREFIX/include/cudf/lists/detail/sorting.hpp - test -f $PREFIX/include/cudf/lists/count_elements.hpp - test -f $PREFIX/include/cudf/lists/explode.hpp From a470368a50e53cdcc9e6e0c71be56bc46626088a Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 28 Apr 2021 11:20:05 -0600 Subject: [PATCH 34/36] Add `template` keyword to template function calls --- cpp/src/lists/interleave_columns.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index 106efd73de3..f2910588392 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -54,7 +54,7 @@ generate_list_offsets_and_validities(table_view const& input, static_assert(sizeof(size_type) == sizeof(int32_t)); auto list_offsets = make_numeric_column( data_type{type_id::INT32}, num_output_lists + 1, mask_state::UNALLOCATED, stream, mr); - auto const d_offsets = list_offsets->mutable_view().begin(); + auto const d_offsets = list_offsets->mutable_view().template begin(); // The array of int8_t to store validities for list elements. auto validities = rmm::device_uvector(has_null_mask ? num_output_lists : 0, stream); @@ -170,7 +170,7 @@ struct interleave_list_entries_fn { { auto const table_dv_ptr = table_device_view::create(input); auto const comp_fn = compute_string_sizes_and_interleave_lists_fn{ - *table_dv_ptr, output_list_offsets.begin(), has_null_mask}; + *table_dv_ptr, output_list_offsets.template begin(), has_null_mask}; if (has_null_mask) { auto [offsets_column, chars_column, null_mask, null_count] = From 2b5c9e81dcd875fb7c8cae1aa75747c601a70d72 Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Wed, 28 Apr 2021 12:32:35 -0600 Subject: [PATCH 35/36] Rewrite comments and rename variables --- cpp/src/lists/interleave_columns.cu | 23 +++++++++++++---------- 1 file changed, 13 insertions(+), 10 deletions(-) diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu index f2910588392..eb205ae8076 100644 --- a/cpp/src/lists/interleave_columns.cu +++ b/cpp/src/lists/interleave_columns.cu @@ -128,26 +128,29 @@ struct compute_string_sizes_and_interleave_lists_fn { auto const str_offsets = str_col.child(strings_column_view::offsets_column_index).template data(); + // The indices of the strings within the source list. + auto const start_str_idx = list_offsets[list_id]; + auto const end_str_idx = list_offsets[list_id + 1]; + // read_idx and write_idx are indices of string elements. size_type write_idx = dst_list_offsets[idx]; - if (not d_chars) { // just compute sizes of strings within a list - for (auto read_idx = list_offsets[list_id], end_idx = list_offsets[list_id + 1]; - read_idx < end_idx; - ++read_idx, ++write_idx) { + + if (not d_chars) { // just compute sizes and validities of strings within a list + for (auto read_idx = start_str_idx; read_idx < end_str_idx; ++read_idx, ++write_idx) { if (has_null_mask) { d_validities[write_idx] = static_cast(str_col.is_valid(read_idx)); } d_offsets[write_idx] = str_offsets[read_idx + 1] - str_offsets[read_idx]; } } else { // just copy the entire memory region containing all strings in the list - // start_idx and end_idx are indices of character elements. - auto const start_idx = str_offsets[list_offsets[list_id]]; - auto const end_idx = str_offsets[list_offsets[list_id + 1]]; - if (start_idx < end_idx) { + // start_byte and end_byte are indices of character of the string elements. + auto const start_byte = str_offsets[start_str_idx]; + auto const end_byte = str_offsets[end_str_idx]; + if (start_byte < end_byte) { auto const input_ptr = - str_col.child(strings_column_view::chars_column_index).template data() + start_idx; + str_col.child(strings_column_view::chars_column_index).template data() + start_byte; auto const output_ptr = d_chars + d_offsets[write_idx]; - thrust::copy(thrust::seq, input_ptr, input_ptr + end_idx - start_idx, output_ptr); + thrust::copy(thrust::seq, input_ptr, input_ptr + end_byte - start_byte, output_ptr); } } } From cca057c818be910fb6245b159dcad56684d54dfc Mon Sep 17 00:00:00 2001 From: Nghia Truong Date: Fri, 30 Apr 2021 16:49:38 -0600 Subject: [PATCH 36/36] Rewrite unit tests --- .../reshape/interleave_columns_tests.cpp | 35 +++++++++++-------- 1 file changed, 21 insertions(+), 14 deletions(-) diff --git a/cpp/tests/reshape/interleave_columns_tests.cpp b/cpp/tests/reshape/interleave_columns_tests.cpp index 5783eeb4526..68cf19bbe9c 100644 --- a/cpp/tests/reshape/interleave_columns_tests.cpp +++ b/cpp/tests/reshape/interleave_columns_tests.cpp @@ -17,6 +17,7 @@ #include #include #include +#include #include #include @@ -374,23 +375,16 @@ using IntListsCol = cudf::test::lists_column_wrapper; using IntCol = cudf::test::fixed_width_column_wrapper; using TView = cudf::table_view; -constexpr bool print_all{true}; // For debugging +constexpr bool print_all{false}; // For debugging constexpr int32_t null{0}; -auto null_at(cudf::size_type idx) -{ - return cudf::detail::make_counting_transform_iterator(0, [idx](auto i) { return i != idx; }); -} -auto null_at(std::vector const& indices) -{ - return cudf::detail::make_counting_transform_iterator(0, [&indices](auto i) { - return std::find(indices.cbegin(), indices.cend(), i) == indices.cend(); - }); -} +auto all_nulls() { return cudf::test::iterator_all_nulls(); } + +auto null_at(cudf::size_type idx) { return cudf::test::iterator_with_null_at(idx); } -auto all_nulls() +auto null_at(std::vector const& indices) { - return cudf::detail::make_counting_transform_iterator(0, [](auto) { return false; }); + return cudf::test::iterator_with_null_at(cudf::host_span{indices}); } } // namespace @@ -424,7 +418,6 @@ TEST_F(ListsColumnsInterleaveTest, InvalidInput) template struct ListsColumnsInterleaveTypedTest : public cudf::test::BaseFixture { }; -#define ListsCol cudf::test::lists_column_wrapper using TypesForTest = cudf::test::Concat; + auto const col = ListsCol{}.release(); auto const results = cudf::interleave_columns(TView{{col->view(), col->view()}}); CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*col, *results, print_all); @@ -440,6 +435,8 @@ TYPED_TEST(ListsColumnsInterleaveTypedTest, InterleaveEmptyColumns) TYPED_TEST(ListsColumnsInterleaveTypedTest, InterleaveOneColumnNotNull) { + using ListsCol = cudf::test::lists_column_wrapper; + auto const col = ListsCol{{1, 2}, {3, 4}, {5, 6}}.release(); auto const results = cudf::interleave_columns(TView{{col->view()}}); CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(*col, *results, print_all); @@ -447,6 +444,8 @@ TYPED_TEST(ListsColumnsInterleaveTypedTest, InterleaveOneColumnNotNull) TYPED_TEST(ListsColumnsInterleaveTypedTest, InterleaveOneColumnWithNulls) { + using ListsCol = cudf::test::lists_column_wrapper; + auto const col = ListsCol{{ListsCol{{1, 2, null}, null_at(2)}, ListsCol{} /*NULL*/, ListsCol{{null, 3, 4, 4, 4, 4}, null_at(0)}, @@ -459,6 +458,8 @@ TYPED_TEST(ListsColumnsInterleaveTypedTest, InterleaveOneColumnWithNulls) TYPED_TEST(ListsColumnsInterleaveTypedTest, SimpleInputNoNull) { + using ListsCol = cudf::test::lists_column_wrapper; + auto const col1 = ListsCol{{1, 2}, {3, 4}, {5, 6}}.release(); auto const col2 = ListsCol{{7, 8}, {9, 10}, {11, 12}}.release(); auto const expected = ListsCol{{1, 2}, {7, 8}, {3, 4}, {9, 10}, {5, 6}, {11, 12}}.release(); @@ -488,6 +489,8 @@ TEST_F(ListsColumnsInterleaveTest, SimpleInputStringsColumnsNoNull) TYPED_TEST(ListsColumnsInterleaveTypedTest, SimpleInputWithNulls) { + using ListsCol = cudf::test::lists_column_wrapper; + auto const col1 = ListsCol{{ListsCol{{1, null, 3, 4}, null_at(1)}, ListsCol{{null, 2, 3, 4}, null_at(0)}, ListsCol{{null, 2, 3, 4}, null_at(0)}, @@ -567,6 +570,8 @@ TEST_F(ListsColumnsInterleaveTest, SimpleInputStringsColumnsWithNulls) TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputNoNull) { + using ListsCol = cudf::test::lists_column_wrapper; + auto const col = ListsCol{{1, 2, 3}, {2, 3}, {3, 4, 5, 6}, {5, 6}, {}, {7}}.release(); auto const col1 = cudf::slice(col->view(), {0, 3})[0]; auto const col2 = cudf::slice(col->view(), {1, 4})[0]; @@ -591,6 +596,8 @@ TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputNoNull) TYPED_TEST(ListsColumnsInterleaveTypedTest, SlicedColumnsInputWithNulls) { + using ListsCol = cudf::test::lists_column_wrapper; + auto const col = ListsCol{{ListsCol{{null, 2, 3}, null_at(0)}, ListsCol{2, 3}, /*NULL*/ ListsCol{{3, null, 5, 6}, null_at(1)},