diff --git a/cpp/include/cudf/detail/scatter.cuh b/cpp/include/cudf/detail/scatter.cuh index 410cd213618..d71a8d0ec24 100644 --- a/cpp/include/cudf/detail/scatter.cuh +++ b/cpp/include/cudf/detail/scatter.cuh @@ -305,7 +305,7 @@ struct column_scatterer_impl { [](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); + scatter_to_gather(scatter_map_begin, scatter_map_end, target.size(), stream); gather_bitmask(cudf::table_view{std::vector{structs_src.child_begin(), structs_src.child_end()}}, gather_map.begin(), diff --git a/cpp/src/copying/copy.cu b/cpp/src/copying/copy.cu index 9f8e6f7bdcb..b0de9cd750e 100644 --- a/cpp/src/copying/copy.cu +++ b/cpp/src/copying/copy.cu @@ -195,21 +195,6 @@ std::unique_ptr scatter_gather_based_if_else(Left const& lhs, { if constexpr (std::is_same::value && std::is_same::value) { - auto const null_map_entry = size + 1; // Out of bounds index, for gather() to nullify. - - auto const gather_lhs = make_counting_transform_iterator( - size_type{0}, lhs_gather_map_functor{is_left, null_map_entry}); - - auto const lhs_gathered_columns = - cudf::detail::gather(table_view{std::vector{lhs}}, - gather_lhs, - gather_lhs + size, - out_of_bounds_policy::NULLIFY, - stream, - mr) - ->release(); - auto& lhs_partial_output = lhs_gathered_columns[0]; - auto scatter_map_rhs = rmm::device_uvector{static_cast(size), stream}; auto const scatter_map_end = thrust::copy_if(rmm::exec_policy(stream), thrust::make_counting_iterator(size_type{0}), @@ -227,7 +212,7 @@ std::unique_ptr scatter_gather_based_if_else(Left const& lhs, table_view{std::vector{scatter_src_rhs->get_column(0).view()}}, scatter_map_rhs.begin(), scatter_map_end, - table_view{std::vector{lhs_partial_output->view()}}, + table_view{std::vector{lhs}}, false, stream, mr); diff --git a/cpp/tests/copying/copy_if_else_nested_tests.cpp b/cpp/tests/copying/copy_if_else_nested_tests.cpp index 9ac34a3044e..7cd56f0ea43 100644 --- a/cpp/tests/copying/copy_if_else_nested_tests.cpp +++ b/cpp/tests/copying/copy_if_else_nested_tests.cpp @@ -102,6 +102,35 @@ TYPED_TEST(TypedCopyIfElseNestedTest, StructsWithNulls) CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(result_column->view(), expected_result->view()); } +TYPED_TEST(TypedCopyIfElseNestedTest, LongerStructsWithNulls) +{ + using T = TypeParam; + + using namespace cudf; + using namespace cudf::test; + + using ints = fixed_width_column_wrapper; + using structs = structs_column_wrapper; + using bools = fixed_width_column_wrapper; + + auto selector_column = bools{1, 1, 1, 1, 0, 1, 0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 0, 1, 0, 1, 0, 0, + 0, 0, 1, 1, 1, 1, 0, 0, 1, 0, 1, 1, 0, 0, 1, 0, 0, 1, 0, 0, 1, 0, + 0, 1, 0, 0, 0, 1, 0, 0, 1, 0, 1, 0, 1, 1, 1, 1, 1, 0, 1, 1, 1, 0} + .release(); + auto lhs_child_1 = + ints{{27, -80, -24, 76, -56, 42, 5, 13, -69, -77, 61, -77, 72, 0, 31, 118, -30, + 86, 125, 0, 0, 0, 75, -49, 125, 60, 116, 118, 64, 20, -70, -18, 0, -25, + 22, -46, -89, -9, 27, -56, -77, 123, 0, -90, 87, -113, -37, 22, -22, -53, 73, + 99, 113, -2, -24, 113, 75, 6, 82, -58, 122, -123, -127, 19, -62, -24}, + iterator_with_null_at(std::vector{13, 19, 20, 21, 32, 42})}; + + auto lhs_structs_column = structs{{lhs_child_1}}.release(); + auto result_column = + copy_if_else(lhs_structs_column->view(), lhs_structs_column->view(), selector_column->view()); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(result_column->view(), lhs_structs_column->view()); +} + TYPED_TEST(TypedCopyIfElseNestedTest, Lists) { using T = TypeParam; diff --git a/cpp/tests/copying/scatter_struct_tests.cpp b/cpp/tests/copying/scatter_struct_tests.cpp index 4ca805f2c18..c07ec7ca7f7 100644 --- a/cpp/tests/copying/scatter_struct_tests.cpp +++ b/cpp/tests/copying/scatter_struct_tests.cpp @@ -241,3 +241,22 @@ TYPED_TEST(TypedStructScatterTest, SourceSmallerThanTargetScatterTest) CUDF_TEST_EXPECT_COLUMNS_EQUAL( *structs_expected, scatter_structs(structs_src, structs_tgt, scatter_map), print_all); } + +TYPED_TEST(TypedStructScatterTest, IntStructNullMaskRegression) +{ + using col_wrapper = cudf::test::fixed_width_column_wrapper; + + auto child_tgt = col_wrapper({0, null, 2}, null_at(1)); + auto struct_col_tgt = structs_col({child_tgt}).release(); + + auto child_src = col_wrapper{20}; + auto struct_col_src = structs_col({child_src}).release(); + + auto scatter_map = int32s_col{2}.release(); + + auto expected_child = col_wrapper({0, null, 20}, null_at(1)); + auto expected_struct = structs_col({expected_child}).release(); + + auto const result = scatter_structs(struct_col_src, struct_col_tgt, scatter_map); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*expected_struct, result, print_all); +}