Fix string functors to adapt to the new make_strings_children from …

…cudf (#2034)

* Fix string functors

Signed-off-by: Nghia Truong <[email protected]>

* Fix style

Signed-off-by: Nghia Truong <[email protected]>

* Include changes from cudf

Signed-off-by: Nghia Truong <[email protected]>

---------

Signed-off-by: Nghia Truong <[email protected]>

src/main/cpp/src/cast_decimal_to_string.cu

@@ -52,8 +52,9 @@ namespace {
 template <typename DecimalType>
 struct decimal_to_non_ansi_string_fn {
   column_device_view d_decimals;
-  size_type* d_offsets{};
-  char* d_chars{};
+  cudf::size_type* d_sizes;
+  char* d_chars;
+  cudf::detail::input_offsetalator d_offsets;
 
   /**
    * @brief Calculates the size of the string required to convert the element, in base-10 format.
@@ -162,13 +163,13 @@ struct decimal_to_non_ansi_string_fn {
   __device__ void operator()(size_type idx)
   {
     if (d_decimals.is_null(idx)) {
-      if (d_chars == nullptr) { d_offsets[idx] = 0; }
+      if (d_chars == nullptr) { d_sizes[idx] = 0; }
       return;
     }
     if (d_chars != nullptr) {
       decimal_to_non_ansi_string(idx);
     } else {
-      d_offsets[idx] = compute_output_size(d_decimals.element<DecimalType>(idx));
+      d_sizes[idx] = compute_output_size(d_decimals.element<DecimalType>(idx));
     }
   }
 };

src/main/cpp/src/cast_float_to_string.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.
@@ -34,8 +34,9 @@ namespace {
 template <typename FloatType>
 struct float_to_string_fn {
   cudf::column_device_view d_floats;
-  cudf::size_type* d_offsets;
+  cudf::size_type* d_sizes;
   char* d_chars;
+  cudf::detail::input_offsetalator d_offsets;
 
   __device__ cudf::size_type compute_output_size(cudf::size_type idx) const
   {
@@ -56,13 +57,13 @@ struct float_to_string_fn {
   __device__ void operator()(cudf::size_type idx) const
   {
     if (d_floats.is_null(idx)) {
-      if (d_chars == nullptr) { d_offsets[idx] = 0; }
+      if (d_chars == nullptr) { d_sizes[idx] = 0; }
       return;
     }
     if (d_chars != nullptr) {
       float_to_string(idx);
     } else {
-      d_offsets[idx] = compute_output_size(idx);
+      d_sizes[idx] = compute_output_size(idx);
     }
   }
 };
@@ -124,4 +125,4 @@ std::unique_ptr<cudf::column> float_to_string(cudf::column_view const& floats,
   return detail::float_to_string(floats, stream, mr);
 }
 
-}  // namespace spark_rapids_jni
+}  // namespace spark_rapids_jni

src/main/cpp/src/format_float.cu

@@ -35,8 +35,9 @@ template <typename FloatType>
 struct format_float_fn {
   cudf::column_device_view d_floats;
   int digits;
-  cudf::size_type* d_offsets;
+  cudf::size_type* d_sizes;
   char* d_chars;
+  cudf::detail::input_offsetalator d_offsets;
 
   __device__ cudf::size_type compute_output_size(FloatType const value) const
   {
@@ -56,13 +57,13 @@ struct format_float_fn {
   __device__ void operator()(cudf::size_type const idx) const
   {
     if (d_floats.is_null(idx)) {
-      if (d_chars == nullptr) { d_offsets[idx] = 0; }
+      if (d_chars == nullptr) { d_sizes[idx] = 0; }
       return;
     }
     if (d_chars != nullptr) {
       format_float(idx);
     } else {
-      d_offsets[idx] = compute_output_size(d_floats.element<FloatType>(idx));
+      d_sizes[idx] = compute_output_size(d_floats.element<FloatType>(idx));
     }
   }
 };
@@ -128,4 +129,4 @@ std::unique_ptr<cudf::column> format_float(cudf::column_view const& floats,
   return detail::format_float(floats, digits, stream, mr);
 }
 
-}  // namespace spark_rapids_jni
+}  // namespace spark_rapids_jni

src/main/cpp/src/map_utils.cu

@@ -16,10 +16,6 @@
 
 #include "map_utils_debug.cuh"
 
-//
-#include <limits>
-
-//
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/null_mask.hpp>
@@ -31,11 +27,11 @@
 #include <cudf/strings/string_view.hpp>
 #include <cudf/strings/strings_column_view.hpp>
 
-//
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 
-//
+#include <cub/device/device_radix_sort.cuh>
+#include <cuda/functional>
 #include <thrust/binary_search.h>
 #include <thrust/copy.h>
 #include <thrust/count.h>
@@ -51,9 +47,7 @@
 #include <thrust/transform.h>
 #include <thrust/transform_reduce.h>
 
-//
-#include <cub/device/device_radix_sort.cuh>
-#include <cuda/functional>
+#include <limits>
 
 namespace spark_rapids_jni {
 
@@ -520,8 +514,9 @@ struct substring_fn {
   cudf::device_span<char const> const d_string;
   cudf::device_span<thrust::pair<SymbolOffsetT, SymbolOffsetT> const> const d_ranges;
 
-  cudf::size_type* d_offsets{};
-  char* d_chars{};
+  cudf::size_type* d_sizes;
+  char* d_chars;
+  cudf::detail::input_offsetalator d_offsets;
 
   __device__ void operator()(cudf::size_type const idx)
   {
@@ -530,7 +525,7 @@ struct substring_fn {
     if (d_chars) {
       memcpy(d_chars + d_offsets[idx], d_string.data() + range.first, size);
     } else {
-      d_offsets[idx] = size;
+      d_sizes[idx] = size;
     }
   }
 };