From 19d8d48df886981a5f4a067cc952827dac1de0ce Mon Sep 17 00:00:00 2001 From: Nghia Truong <7416935+ttnghia@users.noreply.github.com> Date: Mon, 13 May 2024 12:40:25 -0700 Subject: [PATCH] Fix string functors to adapt to the new `make_strings_children` from cudf (#2034) * Fix string functors Signed-off-by: Nghia Truong * Fix style Signed-off-by: Nghia Truong * Include changes from cudf Signed-off-by: Nghia Truong --------- Signed-off-by: Nghia Truong --- src/main/cpp/src/cast_decimal_to_string.cu | 9 +++++---- src/main/cpp/src/cast_float_to_string.cu | 11 ++++++----- src/main/cpp/src/format_float.cu | 9 +++++---- src/main/cpp/src/map_utils.cu | 19 +++++++------------ thirdparty/cudf | 2 +- 5 files changed, 24 insertions(+), 26 deletions(-) diff --git a/src/main/cpp/src/cast_decimal_to_string.cu b/src/main/cpp/src/cast_decimal_to_string.cu index 9d0e27ed59..91b155dae4 100644 --- a/src/main/cpp/src/cast_decimal_to_string.cu +++ b/src/main/cpp/src/cast_decimal_to_string.cu @@ -52,8 +52,9 @@ namespace { template 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(idx)); + d_sizes[idx] = compute_output_size(d_decimals.element(idx)); } } }; diff --git a/src/main/cpp/src/cast_float_to_string.cu b/src/main/cpp/src/cast_float_to_string.cu index 78cedbbf64..b294ca6f1b 100644 --- a/src/main/cpp/src/cast_float_to_string.cu +++ b/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 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 float_to_string(cudf::column_view const& floats, return detail::float_to_string(floats, stream, mr); } -} // namespace spark_rapids_jni \ No newline at end of file +} // namespace spark_rapids_jni diff --git a/src/main/cpp/src/format_float.cu b/src/main/cpp/src/format_float.cu index 1d537595d7..bc3c85bbcc 100644 --- a/src/main/cpp/src/format_float.cu +++ b/src/main/cpp/src/format_float.cu @@ -35,8 +35,9 @@ template 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(idx)); + d_sizes[idx] = compute_output_size(d_floats.element(idx)); } } }; @@ -128,4 +129,4 @@ std::unique_ptr format_float(cudf::column_view const& floats, return detail::format_float(floats, digits, stream, mr); } -} // namespace spark_rapids_jni \ No newline at end of file +} // namespace spark_rapids_jni diff --git a/src/main/cpp/src/map_utils.cu b/src/main/cpp/src/map_utils.cu index 6ae54f4fe9..529ed1d90c 100644 --- a/src/main/cpp/src/map_utils.cu +++ b/src/main/cpp/src/map_utils.cu @@ -16,10 +16,6 @@ #include "map_utils_debug.cuh" -// -#include - -// #include #include #include @@ -31,11 +27,11 @@ #include #include -// #include #include -// +#include +#include #include #include #include @@ -51,9 +47,7 @@ #include #include -// -#include -#include +#include namespace spark_rapids_jni { @@ -520,8 +514,9 @@ struct substring_fn { cudf::device_span const d_string; cudf::device_span 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; } } }; diff --git a/thirdparty/cudf b/thirdparty/cudf index a4cd1d8776..13f028f01a 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit a4cd1d877631e4554c53b57202564398b758324c +Subproject commit 13f028f01ad043b0d24f3e4a28f4267c02806390