Update to new cudf strings where character data is no longer a child column

src/main/cpp/src/cast_string_to_float.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -679,7 +679,7 @@ std::unique_ptr<column> string_to_float(data_type dtype,
         out->mutable_view().null_mask(),
         ansi_mode ? static_cast<ScalarType*>(ansi_count.get())->data() : nullptr,
         static_cast<ScalarType*>(valid_count.get())->data(),
-        string_col.chars().begin<char>(),
+        string_col.chars_begin(stream),
         string_col.offsets().begin<size_type>(),
         string_col.null_mask(),
         num_rows);
@@ -690,7 +690,7 @@ std::unique_ptr<column> string_to_float(data_type dtype,
         out->mutable_view().null_mask(),
         ansi_mode ? static_cast<ScalarType*>(ansi_count.get())->data() : nullptr,
         static_cast<ScalarType*>(valid_count.get())->data(),
-        string_col.chars().begin<char>(),
+        string_col.chars_begin(stream),
         string_col.offsets().begin<size_type>(),
         string_col.null_mask(),
         num_rows);
@@ -714,7 +714,7 @@ std::unique_ptr<column> string_to_float(data_type dtype,
       dest.resize(string_bounds[1] - string_bounds[0]);
 
       cudaMemcpyAsync(dest.data(),
-                      &string_col.chars().data<char const>()[string_bounds[0]],
+                      &string_col.chars_begin(stream)[string_bounds[0]],
                       string_bounds[1] - string_bounds[0],
                       cudaMemcpyDeviceToHost,
                       stream.value());

src/main/cpp/src/map_utils.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -75,32 +75,36 @@ rmm::device_uvector<char> unify_json_strings(cudf::column_view const& input,
   }
 
   auto const d_strings  = cudf::column_device_view::create(input, stream);
-  auto const chars_size = input.child(cudf::strings_column_view::chars_column_index).size();
+  auto const input_scv  = cudf::strings_column_view{input};
+  auto const chars_size = input_scv.chars_size(stream);
   auto const output_size =
     2l +                                            // two extra bracket characters '[' and ']'
     static_cast<int64_t>(chars_size) +
     static_cast<int64_t>(input.size() - 1) +        // append `,` character between input rows
     static_cast<int64_t>(input.null_count()) * 2l;  // replace null with "{}"
+  // TODO: This assertion eventually needs to be removed.
+  // See https://github.com/NVIDIA/spark-rapids-jni/issues/1707
   CUDF_EXPECTS(output_size <= static_cast<int64_t>(std::numeric_limits<cudf::size_type>::max()),
                "The input json column is too large and causes overflow.");
 
   auto const joined_input = cudf::strings::detail::join_strings(
-    cudf::strings_column_view{input},
+    input_scv,
     cudf::string_scalar(","),   // append `,` character between the input rows
     cudf::string_scalar("{}"),  // replacement for null rows
     stream,
     rmm::mr::get_current_device_resource());
-  auto const joined_input_child =
-    joined_input->child(cudf::strings_column_view::chars_column_index);
-  auto const joined_input_size_bytes = joined_input_child.size();
+  auto const joined_input_scv        = cudf::strings_column_view{*joined_input};
+  auto const joined_input_size_bytes = joined_input_scv.chars_size(stream);
+  // TODO: This assertion requires a stream synchronization, may want to remove at some point.
+  // See https://github.com/NVIDIA/spark-rapids-jni/issues/1707
   CUDF_EXPECTS(joined_input_size_bytes + 2 == output_size, "Incorrect output size computation.");
 
   // We want to concatenate 3 strings: "[" + joined_input + "]".
   // For efficiency, let's use memcpy instead of `cudf::strings::detail::concatenate`.
   auto output = rmm::device_uvector<char>(joined_input_size_bytes + 2, stream);
   CUDF_CUDA_TRY(cudaMemsetAsync(output.data(), static_cast<int>('['), 1, stream.value()));
   CUDF_CUDA_TRY(cudaMemcpyAsync(output.data() + 1,
-                                joined_input_child.view().data<char>(),
+                                joined_input_scv.chars_begin(stream),
                                 joined_input_size_bytes,
                                 cudaMemcpyDefault,
                                 stream.value()));

src/main/cpp/src/parse_uri.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -700,19 +700,20 @@ uri_parts __device__ validate_uri(const char* str, int len)
  *
  * @param in_strings Input string column
  * @param chunk Chunk of URI to return
+ * @param base_ptr Pointer to the start of the character data in the strings column
  * @param out_lengths Number of characters in each decode URL
  * @param out_offsets Offsets to the start of the chunks
  * @param out_validity Bitmask of validity data, updated in function
  */
 __global__ void parse_uri_char_counter(column_device_view const in_strings,
                                        URI_chunks chunk,
+                                       char const* const base_ptr,
                                        size_type* const out_lengths,
                                        size_type* const out_offsets,
                                        bitmask_type* out_validity)
 {
   // thread per row
-  auto const tid      = cudf::detail::grid_1d::global_thread_id();
-  auto const base_ptr = in_strings.child(strings_column_view::chars_column_index).data<char>();
+  auto const tid = cudf::detail::grid_1d::global_thread_id();
 
   for (thread_index_type tidx = tid; tidx < in_strings.size();
        tidx += cudf::detail::grid_1d::grid_stride()) {
@@ -778,17 +779,18 @@ __global__ void parse_uri_char_counter(column_device_view const in_strings,
  * @brief Parse protocol and copy from the input string column to the output char buffer.
  *
  * @param in_strings Input string column
+ * @param base_ptr Pointer to the start of the character data in the strings column
  * @param src_offsets Offset value of source strings in in_strings
  * @param offsets Offset value of each string associated with `out_chars`
  * @param out_chars Character buffer for the output string column
  */
 __global__ void parse_uri(column_device_view const in_strings,
+                          char const* const base_ptr,
                           size_type const* const src_offsets,
                           size_type const* const offsets,
                           char* const out_chars)
 {
-  auto const tid      = cudf::detail::grid_1d::global_thread_id();
-  auto const base_ptr = in_strings.child(strings_column_view::chars_column_index).data<char>();
+  auto const tid = cudf::detail::grid_1d::global_thread_id();
 
   for (thread_index_type tidx = tid; tidx < in_strings.size();
        tidx += cudf::detail::grid_1d::grid_stride()) {
@@ -840,6 +842,7 @@ std::unique_ptr<column> parse_uri(strings_column_view const& input,
   parse_uri_char_counter<<<num_threadblocks, threadblock_size, 0, stream.value()>>>(
     *d_strings,
     chunk,
+    input.chars_begin(stream),
     offsets_mutable_view.begin<size_type>(),
     reinterpret_cast<size_type*>(src_offsets.data()),
     reinterpret_cast<bitmask_type*>(null_mask.data()));
@@ -854,23 +857,23 @@ std::unique_ptr<column> parse_uri(strings_column_view const& input,
   // to the host memory
   auto out_chars_bytes = cudf::detail::get_value<size_type>(offsets_view, offset_count - 1, stream);
 
-  // create the chars column
-  auto chars_column = cudf::strings::detail::create_chars_child_column(out_chars_bytes, stream, mr);
-  auto d_out_chars  = chars_column->mutable_view().data<char>();
+  // create the chars buffer
+  auto d_out_chars = rmm::device_buffer(out_chars_bytes, stream, mr);
 
   // copy the characters from the input column to the output column
   parse_uri<<<num_threadblocks, threadblock_size, 0, stream.value()>>>(
     *d_strings,
+    input.chars_begin(stream),
     reinterpret_cast<size_type*>(src_offsets.data()),
     offsets_column->view().begin<size_type>(),
-    d_out_chars);
+    static_cast<char*>(d_out_chars.data()));
 
   auto null_count =
     cudf::null_count(reinterpret_cast<bitmask_type*>(null_mask.data()), 0, strings_count);
 
   return make_strings_column(strings_count,
                              std::move(offsets_column),
-                             std::move(chars_column),
+                             std::move(d_out_chars),
                              null_count,
                              std::move(null_mask));
 }