diff --git a/cpp/include/cudf/detail/gather.cuh b/cpp/include/cudf/detail/gather.cuh index bf488621d52..e4ed6f10e42 100644 --- a/cpp/include/cudf/detail/gather.cuh +++ b/cpp/include/cudf/detail/gather.cuh @@ -466,15 +466,20 @@ struct column_gatherer_impl { mr); }); - gather_bitmask( - // Table view of struct column. - cudf::table_view{ - std::vector{structs_column.child_begin(), structs_column.child_end()}}, - gather_map_begin, - output_struct_members, - nullify_out_of_bounds ? gather_bitmask_op::NULLIFY : gather_bitmask_op::DONT_CHECK, - stream, - mr); + auto const nullable = std::any_of(structs_column.child_begin(), + structs_column.child_end(), + [](auto const& col) { return col.nullable(); }); + if (nullable) { + gather_bitmask( + // Table view of struct column. + cudf::table_view{ + std::vector{structs_column.child_begin(), structs_column.child_end()}}, + gather_map_begin, + output_struct_members, + nullify_out_of_bounds ? gather_bitmask_op::NULLIFY : gather_bitmask_op::DONT_CHECK, + stream, + mr); + } return cudf::make_structs_column( gather_map_size, @@ -652,11 +657,15 @@ std::unique_ptr gather( mr)); } - gather_bitmask_op const op = bounds_policy == out_of_bounds_policy::NULLIFY - ? gather_bitmask_op::NULLIFY - : gather_bitmask_op::DONT_CHECK; - - gather_bitmask(source_table, gather_map_begin, destination_columns, op, stream, mr); + auto const nullable = bounds_policy == out_of_bounds_policy::NULLIFY || + std::any_of(source_table.begin(), source_table.end(), [](auto const& col) { + return col.nullable(); + }); + if (nullable) { + auto const op = bounds_policy == out_of_bounds_policy::NULLIFY ? gather_bitmask_op::NULLIFY + : gather_bitmask_op::DONT_CHECK; + gather_bitmask(source_table, gather_map_begin, destination_columns, op, stream, mr); + } return std::make_unique
(std::move(destination_columns)); } diff --git a/cpp/include/cudf/detail/scatter.cuh b/cpp/include/cudf/detail/scatter.cuh index 30764b9b89f..c1e246122e5 100644 --- a/cpp/include/cudf/detail/scatter.cuh +++ b/cpp/include/cudf/detail/scatter.cuh @@ -25,6 +25,7 @@ #include #include #include +#include #include #include #include @@ -32,6 +33,8 @@ #include #include +#include + namespace cudf { namespace detail { @@ -42,10 +45,9 @@ namespace detail { * function using the PASSTHROUGH op since the resulting map may contain index * values outside the target's range. * - * First, the gather-map is initialized with invalid entries. - * The gather_rows is used since it should always be outside the target size. - * - * Then, the `output[scatter_map[i]] = i`. + * First, the gather-map is initialized with an invalid index. + * The value `numeric_limits::lowest()` is used since it should always be outside the target size. + * Then, `output[scatter_map[i]] = i` for each `i`. * * @tparam MapIterator Iterator type of the input scatter map. * @param scatter_map_begin Beginning of scatter map. @@ -62,11 +64,16 @@ auto scatter_to_gather(MapIterator scatter_map_begin, { using MapValueType = typename thrust::iterator_traits::value_type; - // The gather_map is initialized with gather_rows value to identify pass-through entries - // when calling the gather_bitmask() which applies a pass-through whenever it finds a + // The gather_map is initialized with `numeric_limits::lowest()` value to identify pass-through + // entries when calling the gather_bitmask() which applies a pass-through whenever it finds a // value outside the range of the target column. - // We'll use the gather_rows value for this since it should always be outside the valid range. - auto gather_map = rmm::device_vector(gather_rows, gather_rows); + // We'll use the `numeric_limits::lowest()` value for this since it should always be outside the + // valid range. + auto gather_map = rmm::device_uvector(gather_rows, stream); + thrust::uninitialized_fill(rmm::exec_policy(stream), + gather_map.begin(), + gather_map.end(), + std::numeric_limits::lowest()); // Convert scatter map to a gather map thrust::scatter( @@ -79,6 +86,39 @@ auto scatter_to_gather(MapIterator scatter_map_begin, return gather_map; } +/** + * @brief Create a complement map of `scatter_to_gather` map + * + * The purpose of this map is to create an identity-mapping for the rows that are not + * touched by the `scatter_map`. + * + * The output result of this mapping is firstly initialized as an identity-mapping + * (i.e., `output[i] = i`). Then, for each value `idx` from `scatter_map`, the value `output[idx]` + * is set to `numeric_limits::lowest()`, which is an invalid, out-of-bound index to identify the + * pass-through entries when calling the `gather_bitmask()` function. + * + */ +template +auto scatter_to_gather_complement(MapIterator scatter_map_begin, + MapIterator scatter_map_end, + size_type gather_rows, + rmm::cuda_stream_view stream) +{ + auto gather_map = rmm::device_uvector(gather_rows, stream); + thrust::sequence(rmm::exec_policy(stream), gather_map.begin(), gather_map.end(), 0); + + auto const out_of_bounds_begin = + thrust::make_constant_iterator(std::numeric_limits::lowest()); + auto const out_of_bounds_end = + out_of_bounds_begin + thrust::distance(scatter_map_begin, scatter_map_end); + thrust::scatter(rmm::exec_policy(stream), + out_of_bounds_begin, + out_of_bounds_end, + scatter_map_begin, + gather_map.begin()); + return gather_map; +} + template struct column_scatterer_impl { std::unique_ptr operator()(...) const { CUDF_FAIL("Unsupported type for scatter."); } @@ -214,6 +254,89 @@ struct column_scatterer { } }; +template <> +struct column_scatterer_impl { + template + std::unique_ptr operator()(column_view const& source, + MapItRoot scatter_map_begin, + MapItRoot scatter_map_end, + column_view const& target, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) const + { + CUDF_EXPECTS(source.num_children() == target.num_children(), + "Scatter source and target are not of the same type."); + + auto const scatter_map_size = std::distance(scatter_map_begin, scatter_map_end); + if (scatter_map_size == 0) { return std::make_unique(target, stream, mr); } + + structs_column_view const structs_src(source); + structs_column_view const structs_target(target); + std::vector> output_struct_members(structs_src.num_children()); + + std::transform(structs_src.child_begin(), + structs_src.child_end(), + structs_target.child_begin(), + output_struct_members.begin(), + [&scatter_map_begin, &scatter_map_end, stream, mr](auto const& source_col, + auto const& target_col) { + return type_dispatcher(source_col.type(), + column_scatterer{}, + source_col, + scatter_map_begin, + scatter_map_end, + target_col, + stream, + mr); + }); + + // We still need to call `gather_bitmask` even when the source's children are not nullable, + // as if the target's children have null_masks, those null_masks need to be updated after + // being scattered onto + auto const child_nullable = std::any_of(structs_src.child_begin(), + structs_src.child_end(), + [](auto const& col) { return col.nullable(); }) or + std::any_of(structs_target.child_begin(), + structs_target.child_end(), + [](auto const& col) { return col.nullable(); }); + if (child_nullable) { + auto const gather_map = + scatter_to_gather(scatter_map_begin, scatter_map_end, source.size(), stream); + gather_bitmask(cudf::table_view{std::vector{structs_src.child_begin(), + structs_src.child_end()}}, + gather_map.begin(), + output_struct_members, + gather_bitmask_op::PASSTHROUGH, + stream, + mr); + } + + // Need to put the result column in a vector to call `gather_bitmask` + std::vector> result; + result.emplace_back(cudf::make_structs_column(source.size(), + std::move(output_struct_members), + 0, + rmm::device_buffer{0, stream, mr}, + stream, + mr)); + + // Only gather bitmask from the target column for the rows that have not been scattered onto + // The bitmask from the source column will be gathered at the top level `scatter()` call + if (target.nullable()) { + auto const gather_map = + scatter_to_gather_complement(scatter_map_begin, scatter_map_end, target.size(), stream); + gather_bitmask(table_view{std::vector{target}}, + gather_map.begin(), + result, + gather_bitmask_op::PASSTHROUGH, + stream, + mr); + } + + return std::move(result.front()); + } +}; + /** * @brief Scatters the rows of the source table into a copy of the target table * according to a scatter map. @@ -278,10 +401,8 @@ std::unique_ptr
scatter( // Transform negative indices to index + target size auto updated_scatter_map_begin = thrust::make_transform_iterator(scatter_map_begin, index_converter{target.num_rows()}); - auto updated_scatter_map_end = thrust::make_transform_iterator(scatter_map_end, index_converter{target.num_rows()}); - auto result = std::vector>(target.num_columns()); std::transform(source.begin(), @@ -299,11 +420,16 @@ std::unique_ptr
scatter( mr); }); - auto gather_map = scatter_to_gather( - updated_scatter_map_begin, updated_scatter_map_end, target.num_rows(), stream); - - gather_bitmask(source, gather_map.begin(), result, gather_bitmask_op::PASSTHROUGH, stream, mr); - + // We still need to call `gather_bitmask` even when the source columns are not nullable, + // as if the target has null_mask, that null_mask needs to be updated after scattering + auto const nullable = + std::any_of(source.begin(), source.end(), [](auto const& col) { return col.nullable(); }) or + std::any_of(target.begin(), target.end(), [](auto const& col) { return col.nullable(); }); + if (nullable) { + auto const gather_map = scatter_to_gather( + updated_scatter_map_begin, updated_scatter_map_end, target.num_rows(), stream); + gather_bitmask(source, gather_map.begin(), result, gather_bitmask_op::PASSTHROUGH, stream, mr); + } return std::make_unique
(std::move(result)); } } // namespace detail diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 082f039054e..cbe8dd3caa4 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -201,24 +201,25 @@ ConfigureTest(SORT_TEST ################################################################################################### # - copying tests --------------------------------------------------------------------------------- ConfigureTest(COPYING_TEST - copying/utility_tests.cpp + copying/concatenate_tests.cu + copying/copy_range_tests.cpp + copying/copy_tests.cu + copying/detail_gather_tests.cu + copying/gather_struct_tests.cu copying/gather_tests.cu copying/gather_str_tests.cu copying/gather_list_tests.cu - copying/segmented_gather_list_tests.cpp - copying/gather_struct_tests.cu - copying/detail_gather_tests.cu + copying/get_value_tests.cpp + copying/pack_tests.cu + copying/sample_tests.cpp copying/scatter_tests.cpp copying/scatter_list_tests.cu - copying/copy_range_tests.cpp + copying/scatter_struct_tests.cu + copying/segmented_gather_list_tests.cpp + copying/shift_tests.cpp copying/slice_tests.cpp copying/split_tests.cpp - copying/copy_tests.cu - copying/shift_tests.cpp - copying/get_value_tests.cpp - copying/sample_tests.cpp - copying/concatenate_tests.cu - copying/pack_tests.cu) + copying/utility_tests.cpp) ################################################################################################### # - utilities tests ------------------------------------------------------------------------------- diff --git a/cpp/tests/copying/scatter_struct_tests.cu b/cpp/tests/copying/scatter_struct_tests.cu new file mode 100644 index 00000000000..a9bb1980d53 --- /dev/null +++ b/cpp/tests/copying/scatter_struct_tests.cu @@ -0,0 +1,293 @@ +/* + * Copyright (c) 2020-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 + +using bools_col = cudf::test::fixed_width_column_wrapper; +using int32s_col = cudf::test::fixed_width_column_wrapper; +using structs_col = cudf::test::structs_column_wrapper; +using strings_col = cudf::test::strings_column_wrapper; + +constexpr int32_t null{0}; // Mark for null child elements +constexpr int32_t XXX{0}; // Mark for null struct elements + +template +struct TypedStructScatterTest : public cudf::test::BaseFixture { +}; + +using TestTypes = cudf::test::Concat; + +TYPED_TEST_CASE(TypedStructScatterTest, TestTypes); + +namespace { +void test_scatter(std::unique_ptr const& structs_src, + std::unique_ptr const& structs_tgt, + std::unique_ptr const& structs_expected, + std::unique_ptr const& scatter_map) +{ + auto const source = cudf::table_view{std::vector{structs_src->view()}}; + auto const target = cudf::table_view{std::vector{structs_tgt->view()}}; + auto const result = cudf::scatter(source, scatter_map->view(), target); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(structs_expected->view(), result->get_column(0)); +} +} // namespace + +// Test case when all input columns are empty +TYPED_TEST(TypedStructScatterTest, EmptyInputTest) +{ + using col_wrapper = cudf::test::fixed_width_column_wrapper; + + auto child_col_src = col_wrapper{}; + auto const structs_src = structs_col{{child_col_src}, std::vector{}}.release(); + + auto child_col_tgt = col_wrapper{}; + auto const structs_tgt = structs_col{{child_col_tgt}, std::vector{}}.release(); + + auto const scatter_map = int32s_col{}.release(); + test_scatter(structs_src, structs_tgt, structs_src, scatter_map); + test_scatter(structs_src, structs_tgt, structs_tgt, scatter_map); +} + +// Test case when only the scatter map is empty +TYPED_TEST(TypedStructScatterTest, EmptyScatterMapTest) +{ + using col_wrapper = cudf::test::fixed_width_column_wrapper; + + auto child_col_src = + col_wrapper{{0, 1, 2, 3, null, XXX}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 4; })}; + auto const structs_src = structs_col{ + {child_col_src}, cudf::detail::make_counting_transform_iterator(0, [](auto i) { + return i != 5; + })}.release(); + + auto child_col_tgt = + col_wrapper{{50, null, 70, XXX, 90, 100}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 1; })}; + auto const structs_tgt = structs_col{ + {child_col_tgt}, cudf::detail::make_counting_transform_iterator(0, [](auto i) { + return i != 3; + })}.release(); + + auto const scatter_map = int32s_col{}.release(); + test_scatter(structs_src, structs_tgt, structs_tgt, scatter_map); +} + +TYPED_TEST(TypedStructScatterTest, ScatterAsCopyTest) +{ + using col_wrapper = cudf::test::fixed_width_column_wrapper; + + auto child_col_src = + col_wrapper{{0, 1, 2, 3, null, XXX}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 4; })}; + auto const structs_src = structs_col{ + {child_col_src}, cudf::detail::make_counting_transform_iterator(0, [](auto i) { + return i != 5; + })}.release(); + + auto child_col_tgt = + col_wrapper{{50, null, 70, XXX, 90, 100}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 1; })}; + auto const structs_tgt = structs_col{ + {child_col_tgt}, cudf::detail::make_counting_transform_iterator(0, [](auto i) { + return i != 3; + })}.release(); + + // Scatter as copy: the target should be the same as source + auto const scatter_map = int32s_col{0, 1, 2, 3, 4, 5}.release(); + test_scatter(structs_src, structs_tgt, structs_src, scatter_map); +} + +TYPED_TEST(TypedStructScatterTest, ScatterAsLeftShiftTest) +{ + using col_wrapper = cudf::test::fixed_width_column_wrapper; + + auto child_col_src = + col_wrapper{{0, 1, 2, 3, null, XXX}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 4; })}; + auto const structs_src = structs_col{ + {child_col_src}, cudf::detail::make_counting_transform_iterator(0, [](auto i) { + return i != 5; + })}.release(); + + auto child_col_tgt = + col_wrapper{{50, null, 70, XXX, 90, 100}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 1; })}; + auto const structs_tgt = structs_col{ + {child_col_tgt}, cudf::detail::make_counting_transform_iterator(0, [](auto i) { + return i != 3; + })}.release(); + + auto child_col_expected = + col_wrapper{{2, 3, null, XXX, 0, 1}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 2; })}; + auto structs_expected = structs_col{ + {child_col_expected}, cudf::detail::make_counting_transform_iterator(0, [](auto i) { + return i != 3; + })}.release(); + + auto const scatter_map = int32s_col{-2, -1, 0, 1, 2, 3}.release(); + test_scatter(structs_src, structs_tgt, structs_expected, scatter_map); +} + +TYPED_TEST(TypedStructScatterTest, SimpleScatterTests) +{ + using col_wrapper = cudf::test::fixed_width_column_wrapper; + + // Source data + auto child_col_src = + col_wrapper{{0, 1, 2, 3, null, XXX}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 4; })}; + auto const structs_src = structs_col{ + {child_col_src}, cudf::detail::make_counting_transform_iterator(0, [](auto i) { + return i != 5; + })}.release(); + + // Target data + auto child_col_tgt = + col_wrapper{{50, null, 70, XXX, 90, 100}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 1; })}; + auto const structs_tgt = structs_col{ + {child_col_tgt}, cudf::detail::make_counting_transform_iterator(0, [](auto i) { + return i != 3; + })}.release(); + + // Expected data + auto child_col_expected1 = + col_wrapper{{1, null, 70, XXX, 0, 2}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 1; })}; + auto const structs_expected1 = structs_col{ + {child_col_expected1}, cudf::detail::make_counting_transform_iterator(0, [](auto i) { + return i != 3; + })}.release(); + auto const scatter_map1 = int32s_col{-2, 0, 5}.release(); + test_scatter(structs_src, structs_tgt, structs_expected1, scatter_map1); + + // Expected data + auto child_col_expected2 = + col_wrapper{{1, null, 70, 3, 0, 2}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 1; })}; + auto const structs_expected2 = structs_col{ + {child_col_expected2}, cudf::detail::make_counting_transform_iterator(0, [](auto i) { + return true; + })}.release(); + auto const scatter_map2 = int32s_col{-2, 0, 5, 3}.release(); + test_scatter(structs_src, structs_tgt, structs_expected2, scatter_map2); +} + +TYPED_TEST(TypedStructScatterTest, ComplexDataScatterTest) +{ + // Testing scatter() on struct. + using col_wrapper = cudf::test::fixed_width_column_wrapper; + + // Source data + auto names_column_src = + strings_col{{"Newton", "Washington", "Cherry", "Kiwi", "Lemon", "Tomato"}, + cudf::detail::make_counting_transform_iterator(0, [](auto) { return true; })}; + auto ages_column_src = + col_wrapper{{5, 10, 15, 20, 25, 30}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 4; })}; + auto is_human_col_src = + bools_col{{true, true, false, false, false, false}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 3; })}; + + // Target data + auto names_column_tgt = + strings_col{{"String 0", "String 1", "String 2", "String 3", "String 4", "String 5"}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0; })}; + auto ages_column_tgt = + col_wrapper{{50, 60, 70, 80, 90, 100}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 1; })}; + auto is_human_col_tgt = + bools_col{{true, true, true, true, true, true}, + cudf::detail::make_counting_transform_iterator(0, [](auto) { return true; })}; + + // Expected data + auto names_column_expected = + strings_col{{"String 0", "Lemon", "Kiwi", "Cherry", "Washington", "Newton"}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0; })}; + auto ages_column_expected = + col_wrapper{{50, 25, 20, 15, 10, 5}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 1; })}; + auto is_human_col_expected = + bools_col{{true, false, false, false, true, true}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 2; })}; + + auto const structs_src = structs_col{ + {names_column_src, ages_column_src, is_human_col_src}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { + return i != 5; + })}.release(); + auto const structs_tgt = structs_col{ + {names_column_tgt, ages_column_tgt, is_human_col_tgt}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { + return i != 2; + })}.release(); + auto const structs_expected = structs_col{ + {names_column_expected, ages_column_expected, is_human_col_expected}, + cudf::detail::make_counting_transform_iterator(0, [](auto i) { + return true; + })}.release(); + + // The first element of the target is not overwritten + auto const scatter_map = int32s_col{-1, 4, 3, 2, 1}.release(); + test_scatter(structs_src, structs_tgt, structs_expected, scatter_map); +} + +TYPED_TEST(TypedStructScatterTest, ScatterStructOfListsTest) +{ + // Testing gather() on struct> + using lists_col = cudf::test::lists_column_wrapper; + + // Source data + auto lists_col_src = + lists_col{{{5}, {10, 15}, {20, 25, 30}, {35, 40, 45, 50}, {55, 60, 65}, {70, 75}, {80}, {}, {}}, + // Valid for elements 0, 3, 6,... + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return !(i % 3); })}; + auto const structs_src = structs_col{{lists_col_src}}.release(); + + // Target data + auto lists_col_tgt = + lists_col{{{1}, {2, 3}, {4, 5, 6}, {7, 8}, {9}, {10, 11, 12, 13}, {}, {14}, {15, 16}}, + // Valid for elements 1, 3, 5, 7,... + cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i % 2; })}; + auto const structs_tgt = structs_col{{lists_col_tgt}}.release(); + + // Expected data + auto const validity_expected = std::vector{0, 1, 1, 0, 0, 1, 1, 0, 0}; + auto lists_col_expected = lists_col{ + {{1}, {2, 3}, {80}, {70, 75}, {55, 60, 65}, {35, 40, 45, 50}, {5}, {10, 15}, {20, 25, 30}}, + validity_expected.begin()}; + auto const structs_expected = structs_col{{lists_col_expected}}.release(); + + // The first 2 elements of the target is not overwritten + auto const scatter_map = int32s_col{-3, -2, -1, 5, 4, 3, 2}.release(); + test_scatter(structs_src, structs_tgt, structs_expected, scatter_map); +} diff --git a/cpp/tests/partitioning/partition_test.cpp b/cpp/tests/partitioning/partition_test.cpp index a6838112a54..ed994da20f8 100644 --- a/cpp/tests/partitioning/partition_test.cpp +++ b/cpp/tests/partitioning/partition_test.cpp @@ -141,6 +141,35 @@ TYPED_TEST(PartitionTest, Identity) run_partition_test(table_to_partition, map, 6, table_to_partition, expected_offsets); } +TYPED_TEST(PartitionTest, Struct) +{ + using value_type = cudf::test::GetType; + using map_type = cudf::test::GetType; + + fixed_width_column_wrapper A({1, 2}, {0, 1}); + auto struct_col = cudf::test::structs_column_wrapper({A}, {0, 1}).release(); + auto table_to_partition = cudf::table_view{{*struct_col}}; + + fixed_width_column_wrapper map{9, 2}; + + fixed_width_column_wrapper A_expected({2, 1}, {1, 0}); + auto struct_expected = cudf::test::structs_column_wrapper({A_expected}, {1, 0}).release(); + auto expected = cudf::table_view{{*struct_expected}}; + + std::vector expected_offsets{0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 2, 2, 2}; + + // This does not work because we cannot sort a struct right now... + // run_partition_test(table_to_partition, map, 12, expected, expected_offsets); + // But there is no ambiguity in the ordering so I'll just copy it all here for now. + auto num_partitions = 12; + auto result = cudf::partition(table_to_partition, map, num_partitions); + auto const& actual_partitioned_table = result.first; + auto const& actual_offsets = result.second; + EXPECT_EQ(actual_offsets, expected_offsets); + + CUDF_TEST_EXPECT_TABLES_EQUAL(expected, *actual_partitioned_table); +} + TYPED_TEST(PartitionTest, Reverse) { using value_type = cudf::test::GetType;