Additional fixes for libcudf large strings support

cpp/src/io/json/host_tree_algorithms.cu

@@ -9,6 +9,7 @@
 #include <cudf/detail/algorithms/copy_if.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/detail/offsets_iterator_factory.cuh>
 #include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/detail/utilities/visitor_overload.hpp>
 #include <cudf/strings/strings_column_view.hpp>
@@ -26,6 +27,7 @@
 #include <cuda/iterator>
 #include <cuda/std/iterator>
 #include <cuda/std/tuple>
+#include <thrust/copy.h>
 #include <thrust/for_each.h>
 #include <thrust/scan.h>
 #include <thrust/scatter.h>
@@ -120,10 +122,15 @@ std::vector<std::string> copy_strings_to_host_sync(
     auto const scv     = cudf::strings_column_view(col);
     auto const h_chars = cudf::detail::make_host_vector_async<char>(
       cudf::device_span<char const>(scv.chars_begin(stream), scv.chars_size(stream)), stream);
+    auto d_offsets = rmm::device_uvector<int64_t>(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<cudf::size_type const>(scv.offsets().data<cudf::size_type>() + scv.offset(),
-                                               scv.size() + 1),
-      stream);
+      cudf::device_span<int64_t const>(d_offsets.data(), d_offsets.size()), stream);
     stream.synchronize();
 
     // build std::string vector from chars and offsets

cpp/src/io/json/json_tree.cu

@@ -10,6 +10,7 @@
 #include <cudf/detail/algorithms/reduce.cuh>
 #include <cudf/detail/cuco_helpers.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/detail/offsets_iterator_factory.cuh>
 #include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/hashing/detail/default_hash.cuh>
 #include <cudf/hashing/detail/hashing.hpp>
@@ -560,23 +561,24 @@ std::pair<size_t, rmm::device_uvector<size_type>> 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<size_type>();
+  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<cudf::string_view>::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<cudf::string_view>{}(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;
   };
 

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 <cuda/std/utility>
 #include <thrust/for_each.h>
+#include <thrust/iterator/constant_iterator.h>
 
 namespace cudf {
 namespace strings {
@@ -157,13 +158,9 @@ std::unique_ptr<column> 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<size_type>(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<column>(std::move(offsets), rmm::device_buffer{}, 0);
-  }();
+  auto sizes_itr      = thrust::constant_iterator(static_cast<size_type>(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 =

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<cudf::column> 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{});