From e0602a8f2f8e79f8b35ad1d5650c2ea7bc6985cb Mon Sep 17 00:00:00 2001 From: David Wendt Date: Mon, 10 Aug 2026 16:43:52 -0400 Subject: [PATCH 1/2] Additional fixes for libcudf large strings support --- cpp/src/io/json/host_tree_algorithms.cu | 13 ++++++++--- cpp/src/io/json/json_tree.cu | 24 ++++++++++---------- cpp/src/strings/combine/join.cu | 13 +++++------ cpp/src/text/ngrams_tokenize.cu | 6 ++--- cpp/tests/groupby/collect_list_tests.cpp | 28 +++++++----------------- 5 files changed, 39 insertions(+), 45 deletions(-) diff --git a/cpp/src/io/json/host_tree_algorithms.cu b/cpp/src/io/json/host_tree_algorithms.cu index cfa616f6ec45..6657672d828e 100644 --- a/cpp/src/io/json/host_tree_algorithms.cu +++ b/cpp/src/io/json/host_tree_algorithms.cu @@ -9,6 +9,7 @@ #include #include #include +#include #include #include #include @@ -26,6 +27,7 @@ #include #include #include +#include #include #include #include @@ -120,10 +122,15 @@ std::vector copy_strings_to_host_sync( auto const scv = cudf::strings_column_view(col); auto const h_chars = cudf::detail::make_host_vector_async( cudf::device_span(scv.chars_begin(stream), scv.chars_size(stream)), stream); + auto d_offsets = rmm::device_uvector(scv.size() + 1, stream); + auto offset_itr = + cudf::detail::offsetalator_factory::make_input_iterator(scv.offsets(), scv.offset()); + thrust::copy(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + offset_itr, + offset_itr + scv.size() + 1, + d_offsets.begin()); auto const h_offsets = cudf::detail::make_host_vector_async( - cudf::device_span(scv.offsets().data() + scv.offset(), - scv.size() + 1), - stream); + cudf::device_span(d_offsets.data(), d_offsets.size()), stream); stream.synchronize(); // build std::string vector from chars and offsets diff --git a/cpp/src/io/json/json_tree.cu b/cpp/src/io/json/json_tree.cu index eaafbcab7bc0..043c6feb2f9f 100644 --- a/cpp/src/io/json/json_tree.cu +++ b/cpp/src/io/json/json_tree.cu @@ -10,6 +10,7 @@ #include #include #include +#include #include #include #include @@ -560,23 +561,24 @@ std::pair> remapped_field_nodes_after_uni // insert and find. -> array // store to static_map with keys as field key[index], and values as key[array[index]] - auto str_view = strings_column_view{utf8_decoded_fields->view()}; - auto const char_ptr = str_view.chars_begin(stream); - auto const offset_ptr = str_view.offsets().begin(); + auto str_view = strings_column_view{utf8_decoded_fields->view()}; + auto const char_ptr = str_view.chars_begin(stream); + auto const offset_itr = + cudf::detail::offsetalator_factory::make_input_iterator(str_view.offsets()); // String hasher auto const d_hasher = cuda::proclaim_return_type< typename cudf::hashing::detail::default_hash::result_type>( - [char_ptr, offset_ptr] __device__(auto node_id) { - auto const field_name = cudf::string_view(char_ptr + offset_ptr[node_id], - offset_ptr[node_id + 1] - offset_ptr[node_id]); + [char_ptr, offset_itr] __device__(auto node_id) { + auto const field_name = cudf::string_view(char_ptr + offset_itr[node_id], + offset_itr[node_id + 1] - offset_itr[node_id]); return cudf::hashing::detail::default_hash{}(field_name); }); - auto const d_equal = [char_ptr, offset_ptr] __device__(auto node_id1, auto node_id2) { - auto const field_name1 = cudf::string_view(char_ptr + offset_ptr[node_id1], - offset_ptr[node_id1 + 1] - offset_ptr[node_id1]); - auto const field_name2 = cudf::string_view(char_ptr + offset_ptr[node_id2], - offset_ptr[node_id2 + 1] - offset_ptr[node_id2]); + auto const d_equal = [char_ptr, offset_itr] __device__(auto node_id1, auto node_id2) { + auto const field_name1 = cudf::string_view(char_ptr + offset_itr[node_id1], + offset_itr[node_id1 + 1] - offset_itr[node_id1]); + auto const field_name2 = cudf::string_view(char_ptr + offset_itr[node_id2], + offset_itr[node_id2 + 1] - offset_itr[node_id2]); return field_name1 == field_name2; }; diff --git a/cpp/src/strings/combine/join.cu b/cpp/src/strings/combine/join.cu index aadd3cc83752..6534858eb10c 100644 --- a/cpp/src/strings/combine/join.cu +++ b/cpp/src/strings/combine/join.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -26,6 +26,7 @@ #include #include +#include namespace cudf { namespace strings { @@ -157,13 +158,9 @@ std::unique_ptr join_strings(strings_column_view const& input, std::overflow_error); // build the offsets: single string output has offsets [0,chars-size] - auto offsets_column = [&] { - auto h_offsets = cudf::detail::make_host_vector(2, stream); - h_offsets[0] = 0; - h_offsets[1] = chars.size(); - auto offsets = cudf::detail::make_device_uvector_async(h_offsets, stream, mr); - return std::make_unique(std::move(offsets), rmm::device_buffer{}, 0); - }(); + auto sizes_itr = thrust::constant_iterator(static_cast(chars.size())); + auto offsets_column = std::get<0>( + cudf::strings::detail::make_offsets_child_column(sizes_itr, sizes_itr + 1, stream, mr)); // build the null mask: only one output row so it is either all-valid or all-null auto const null_count = diff --git a/cpp/src/text/ngrams_tokenize.cu b/cpp/src/text/ngrams_tokenize.cu index b07ac676fa77..22d0a1f5887a 100644 --- a/cpp/src/text/ngrams_tokenize.cu +++ b/cpp/src/text/ngrams_tokenize.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -221,8 +221,8 @@ std::unique_ptr ngrams_tokenize(cudf::strings_column_view const& s d_ngram_offsets, ngram_sizes.data()}); // build the offsets column -- converting the ngram sizes into offsets - auto offsets_column = std::get<0>( - cudf::detail::make_offsets_child_column(ngram_sizes.begin(), ngram_sizes.end(), stream, mr)); + auto offsets_column = std::get<0>(cudf::strings::detail::make_offsets_child_column( + ngram_sizes.begin(), ngram_sizes.end(), stream, mr)); // create the output strings column return make_strings_column( total_ngrams, std::move(offsets_column), chars.release(), 0, rmm::device_buffer{}); diff --git a/cpp/tests/groupby/collect_list_tests.cpp b/cpp/tests/groupby/collect_list_tests.cpp index b416b22b6a86..c4e00b742796 100644 --- a/cpp/tests/groupby/collect_list_tests.cpp +++ b/cpp/tests/groupby/collect_list_tests.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -134,7 +134,7 @@ TYPED_TEST(groupby_collect_list_test, CollectOnEmptyInputLists) using LCW = cudf::test::lists_column_wrapper; - auto offsets = cudf::data_type{cudf::type_to_id()}; + auto offsets = cudf::data_type{cudf::type_id::INT32}; cudf::test::fixed_width_column_wrapper keys{}; auto values = @@ -162,30 +162,18 @@ TYPED_TEST(groupby_collect_list_test, CollectOnEmptyInputListsOfStructs) auto struct_child = LCW{}; auto struct_column = cudf::test::structs_column_wrapper{{struct_child}}; - auto values = - cudf::make_lists_column(0, - cudf::make_empty_column(cudf::type_to_id()), - struct_column.release(), - 0, - {}); + auto values = cudf::make_lists_column( + 0, cudf::make_empty_column(cudf::type_id::INT32), struct_column.release(), 0, {}); cudf::test::fixed_width_column_wrapper expect_keys{}; auto expect_struct_child = LCW{}; auto expect_struct_column = cudf::test::structs_column_wrapper{{expect_struct_child}}; - auto expect_child = - cudf::make_lists_column(0, - cudf::make_empty_column(cudf::type_to_id()), - expect_struct_column.release(), - 0, - {}); - auto expect_values = - cudf::make_lists_column(0, - cudf::make_empty_column(cudf::type_to_id()), - std::move(expect_child), - 0, - {}); + auto expect_child = cudf::make_lists_column( + 0, cudf::make_empty_column(cudf::type_id::INT32), expect_struct_column.release(), 0, {}); + auto expect_values = cudf::make_lists_column( + 0, cudf::make_empty_column(cudf::type_id::INT32), std::move(expect_child), 0, {}); auto agg = cudf::make_collect_list_aggregation(); test_single_agg(keys, values->view(), expect_keys, expect_values->view(), std::move(agg)); From 074a471b736ed11a5bbae1fb02c189597842318c Mon Sep 17 00:00:00 2001 From: David Wendt Date: Mon, 10 Aug 2026 16:48:05 -0400 Subject: [PATCH 2/2] remove unrelated change --- cpp/tests/groupby/collect_list_tests.cpp | 28 +++++++++++++++++------- 1 file changed, 20 insertions(+), 8 deletions(-) diff --git a/cpp/tests/groupby/collect_list_tests.cpp b/cpp/tests/groupby/collect_list_tests.cpp index c4e00b742796..b416b22b6a86 100644 --- a/cpp/tests/groupby/collect_list_tests.cpp +++ b/cpp/tests/groupby/collect_list_tests.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2021-2024, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -134,7 +134,7 @@ TYPED_TEST(groupby_collect_list_test, CollectOnEmptyInputLists) using LCW = cudf::test::lists_column_wrapper; - auto offsets = cudf::data_type{cudf::type_id::INT32}; + auto offsets = cudf::data_type{cudf::type_to_id()}; cudf::test::fixed_width_column_wrapper keys{}; auto values = @@ -162,18 +162,30 @@ TYPED_TEST(groupby_collect_list_test, CollectOnEmptyInputListsOfStructs) auto struct_child = LCW{}; auto struct_column = cudf::test::structs_column_wrapper{{struct_child}}; - auto values = cudf::make_lists_column( - 0, cudf::make_empty_column(cudf::type_id::INT32), struct_column.release(), 0, {}); + auto values = + cudf::make_lists_column(0, + cudf::make_empty_column(cudf::type_to_id()), + struct_column.release(), + 0, + {}); cudf::test::fixed_width_column_wrapper expect_keys{}; auto expect_struct_child = LCW{}; auto expect_struct_column = cudf::test::structs_column_wrapper{{expect_struct_child}}; - auto expect_child = cudf::make_lists_column( - 0, cudf::make_empty_column(cudf::type_id::INT32), expect_struct_column.release(), 0, {}); - auto expect_values = cudf::make_lists_column( - 0, cudf::make_empty_column(cudf::type_id::INT32), std::move(expect_child), 0, {}); + auto expect_child = + cudf::make_lists_column(0, + cudf::make_empty_column(cudf::type_to_id()), + expect_struct_column.release(), + 0, + {}); + auto expect_values = + cudf::make_lists_column(0, + cudf::make_empty_column(cudf::type_to_id()), + std::move(expect_child), + 0, + {}); auto agg = cudf::make_collect_list_aggregation(); test_single_agg(keys, values->view(), expect_keys, expect_values->view(), std::move(agg));