diff --git a/cpp/src/strings/convert/convert_booleans.cu b/cpp/src/strings/convert/convert_booleans.cu index bf73800ad06..6b64006fa24 100644 --- a/cpp/src/strings/convert/convert_booleans.cu +++ b/cpp/src/strings/convert/convert_booleans.cu @@ -16,23 +16,19 @@ #include #include -#include #include #include #include #include -#include +#include #include #include #include -#include -#include #include #include #include -#include #include #include @@ -99,13 +95,14 @@ struct from_booleans_fn { column_device_view const d_column; string_view d_true; string_view d_false; - size_type* d_offsets{}; + size_type* d_sizes{}; char* d_chars{}; + cudf::detail::input_offsetalator d_offsets; __device__ void operator()(size_type idx) const { if (d_column.is_null(idx)) { - if (d_chars == nullptr) { d_offsets[idx] = 0; } + if (d_chars == nullptr) { d_sizes[idx] = 0; } return; } @@ -113,7 +110,7 @@ struct from_booleans_fn { auto const result = d_column.element(idx) ? d_true : d_false; memcpy(d_chars + d_offsets[idx], result.data(), result.size_bytes()); } else { - d_offsets[idx] = d_column.element(idx) ? d_true.size_bytes() : d_false.size_bytes(); + d_sizes[idx] = d_column.element(idx) ? d_true.size_bytes() : d_false.size_bytes(); } }; }; @@ -143,8 +140,8 @@ std::unique_ptr from_booleans(column_view const& booleans, // copy null mask rmm::device_buffer null_mask = cudf::detail::copy_bitmask(booleans, stream, mr); - auto [offsets, chars] = - make_strings_children(from_booleans_fn{d_column, d_true, d_false}, strings_count, stream, mr); + auto [offsets, chars] = experimental::make_strings_children( + from_booleans_fn{d_column, d_true, d_false}, strings_count, stream, mr); return make_strings_column(strings_count, std::move(offsets), diff --git a/cpp/src/strings/convert/convert_datetime.cu b/cpp/src/strings/convert/convert_datetime.cu index d6449fbb6c8..ddf68eae951 100644 --- a/cpp/src/strings/convert/convert_datetime.cu +++ b/cpp/src/strings/convert/convert_datetime.cu @@ -22,7 +22,7 @@ #include #include #include -#include +#include #include #include #include @@ -37,7 +37,6 @@ #include #include -#include #include #include #include @@ -756,8 +755,9 @@ struct datetime_formatter_fn { column_device_view const d_timestamps; column_device_view const d_format_names; device_span const d_format_items; - size_type* d_offsets{}; + size_type* d_sizes{}; char* d_chars{}; + cudf::detail::input_offsetalator d_offsets; /** * @brief Specialized modulo expression that handles negative values. @@ -1087,14 +1087,14 @@ struct datetime_formatter_fn { __device__ void operator()(size_type idx) const { if (d_timestamps.is_null(idx)) { - if (!d_chars) { d_offsets[idx] = 0; } + if (!d_chars) { d_sizes[idx] = 0; } return; } auto const tstamp = d_timestamps.element(idx); if (d_chars) { timestamp_to_string(tstamp, d_chars + d_offsets[idx]); } else { - d_offsets[idx] = compute_output_size(tstamp); + d_sizes[idx] = compute_output_size(tstamp); } } }; @@ -1109,7 +1109,7 @@ struct dispatch_from_timestamps_fn { rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) const { - return make_strings_children( + return experimental::make_strings_children( datetime_formatter_fn{d_timestamps, d_format_names, d_format_items}, d_timestamps.size(), stream, diff --git a/cpp/src/strings/convert/convert_durations.cu b/cpp/src/strings/convert/convert_durations.cu index 77c750848cf..faf9a83f016 100644 --- a/cpp/src/strings/convert/convert_durations.cu +++ b/cpp/src/strings/convert/convert_durations.cu @@ -17,7 +17,7 @@ #include #include #include -#include +#include #include #include @@ -26,10 +26,8 @@ #include #include -#include #include #include -#include #include #include @@ -192,8 +190,9 @@ struct from_durations_fn { column_device_view d_durations; format_item const* d_format_items; size_type items_count; - size_type* d_offsets{}; + size_type* d_sizes{}; char* d_chars{}; + cudf::detail::input_offsetalator d_offsets; __device__ int8_t format_length(char format_char, duration_component const* const timeparts) const { @@ -378,14 +377,14 @@ struct from_durations_fn { __device__ void operator()(size_type idx) { if (d_durations.is_null(idx)) { - if (d_chars == nullptr) { d_offsets[idx] = 0; } + if (d_chars == nullptr) { d_sizes[idx] = 0; } return; } if (d_chars != nullptr) { set_chars(idx); } else { - d_offsets[idx] = string_size(d_durations.template element(idx)); + d_sizes[idx] = string_size(d_durations.template element(idx)); } } }; @@ -415,11 +414,11 @@ struct dispatch_from_durations_fn { // copy null mask rmm::device_buffer null_mask = cudf::detail::copy_bitmask(durations, stream, mr); - auto [offsets, chars] = - make_strings_children(from_durations_fn{d_column, d_format_items, compiler.items_count()}, - strings_count, - stream, - mr); + auto [offsets, chars] = experimental::make_strings_children( + from_durations_fn{d_column, d_format_items, compiler.items_count()}, + strings_count, + stream, + mr); return make_strings_column(strings_count, std::move(offsets), diff --git a/cpp/src/strings/convert/convert_fixed_point.cu b/cpp/src/strings/convert/convert_fixed_point.cu index 446baa8dea9..34f81b8b407 100644 --- a/cpp/src/strings/convert/convert_fixed_point.cu +++ b/cpp/src/strings/convert/convert_fixed_point.cu @@ -23,7 +23,7 @@ #include #include #include -#include +#include #include #include #include @@ -37,10 +37,7 @@ #include #include #include -#include -#include #include -#include #include namespace cudf { @@ -198,8 +195,9 @@ namespace { template struct from_fixed_point_fn { column_device_view d_decimals; - size_type* d_offsets{}; + size_type* d_sizes{}; char* d_chars{}; + cudf::detail::input_offsetalator d_offsets; /** * @brief Converts a decimal element into a string. @@ -219,13 +217,13 @@ struct from_fixed_point_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) { fixed_point_element_to_string(idx); } else { - d_offsets[idx] = + d_sizes[idx] = fixed_point_string_size(d_decimals.element(idx), d_decimals.type().scale()); } } @@ -244,8 +242,8 @@ struct dispatch_from_fixed_point_fn { auto const d_column = column_device_view::create(input, stream); - auto [offsets, chars] = - make_strings_children(from_fixed_point_fn{*d_column}, input.size(), stream, mr); + auto [offsets, chars] = experimental::make_strings_children( + from_fixed_point_fn{*d_column}, input.size(), stream, mr); return make_strings_column(input.size(), std::move(offsets), diff --git a/cpp/src/strings/convert/convert_floats.cu b/cpp/src/strings/convert/convert_floats.cu index c6061f7d8e6..0ed80b976fd 100644 --- a/cpp/src/strings/convert/convert_floats.cu +++ b/cpp/src/strings/convert/convert_floats.cu @@ -21,7 +21,7 @@ #include #include #include -#include +#include #include #include #include @@ -32,9 +32,7 @@ #include #include -#include #include -#include #include #include @@ -356,8 +354,9 @@ struct ftos_converter { template struct from_floats_fn { column_device_view d_floats; - size_type* d_offsets; + size_type* d_sizes; char* d_chars; + cudf::detail::input_offsetalator d_offsets; __device__ size_type compute_output_size(FloatType value) { @@ -375,13 +374,13 @@ struct from_floats_fn { __device__ void operator()(size_type idx) { 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(d_floats.element(idx)); + d_sizes[idx] = compute_output_size(d_floats.element(idx)); } } }; @@ -404,8 +403,8 @@ struct dispatch_from_floats_fn { // copy the null mask rmm::device_buffer null_mask = cudf::detail::copy_bitmask(floats, stream, mr); - auto [offsets, chars] = - make_strings_children(from_floats_fn{d_column}, strings_count, stream, mr); + auto [offsets, chars] = experimental::make_strings_children( + from_floats_fn{d_column}, strings_count, stream, mr); return make_strings_column(strings_count, std::move(offsets), diff --git a/cpp/src/strings/convert/convert_hex.cu b/cpp/src/strings/convert/convert_hex.cu index 95af378fc3f..1f9fc3858f8 100644 --- a/cpp/src/strings/convert/convert_hex.cu +++ b/cpp/src/strings/convert/convert_hex.cu @@ -19,7 +19,7 @@ #include #include #include -#include +#include #include #include #include @@ -123,8 +123,9 @@ struct dispatch_hex_to_integers_fn { template struct integer_to_hex_fn { column_device_view const d_column; - size_type* d_offsets{}; + size_type* d_sizes{}; char* d_chars{}; + cudf::detail::input_offsetalator d_offsets; __device__ void byte_to_hex(uint8_t byte, char* hex) { @@ -141,7 +142,7 @@ struct integer_to_hex_fn { __device__ void operator()(size_type idx) { if (d_column.is_null(idx)) { - if (!d_chars) { d_offsets[idx] = 0; } + if (!d_chars) { d_sizes[idx] = 0; } return; } @@ -167,7 +168,7 @@ struct integer_to_hex_fn { --byte_index; } } else { - d_offsets[idx] = static_cast(bytes) * 2; // 2 hex characters per byte + d_sizes[idx] = static_cast(bytes) * 2; // 2 hex characters per byte } } }; @@ -181,7 +182,7 @@ struct dispatch_integers_to_hex_fn { { auto const d_column = column_device_view::create(input, stream); - auto [offsets_column, chars] = cudf::strings::detail::make_strings_children( + auto [offsets_column, chars] = experimental::make_strings_children( integer_to_hex_fn{*d_column}, input.size(), stream, mr); return make_strings_column(input.size(), diff --git a/cpp/src/strings/convert/convert_integers.cu b/cpp/src/strings/convert/convert_integers.cu index f3e639817a6..918369ead4d 100644 --- a/cpp/src/strings/convert/convert_integers.cu +++ b/cpp/src/strings/convert/convert_integers.cu @@ -23,7 +23,7 @@ #include #include #include -#include +#include #include #include #include @@ -34,9 +34,7 @@ #include #include -#include #include -#include #include #include #include @@ -314,8 +312,9 @@ namespace { template struct from_integers_fn { column_device_view d_integers; - size_type* d_offsets; + size_type* d_sizes; char* d_chars; + cudf::detail::input_offsetalator d_offsets; /** * @brief Converts an integer element into a string. @@ -334,13 +333,13 @@ struct from_integers_fn { __device__ void operator()(size_type idx) { if (d_integers.is_null(idx)) { - if (d_chars == nullptr) { d_offsets[idx] = 0; } + if (d_chars == nullptr) { d_sizes[idx] = 0; } return; } if (d_chars != nullptr) { integer_element_to_string(idx); } else { - d_offsets[idx] = count_digits(d_integers.element(idx)); + d_sizes[idx] = count_digits(d_integers.element(idx)); } } }; @@ -363,8 +362,8 @@ struct dispatch_from_integers_fn { // copy the null mask rmm::device_buffer null_mask = cudf::detail::copy_bitmask(integers, stream, mr); - auto [offsets, chars] = - make_strings_children(from_integers_fn{d_column}, strings_count, stream, mr); + auto [offsets, chars] = experimental::make_strings_children( + from_integers_fn{d_column}, strings_count, stream, mr); return make_strings_column(strings_count, std::move(offsets), diff --git a/cpp/src/strings/convert/convert_ipv4.cu b/cpp/src/strings/convert/convert_ipv4.cu index 3d259f0ab82..33f6c553001 100644 --- a/cpp/src/strings/convert/convert_ipv4.cu +++ b/cpp/src/strings/convert/convert_ipv4.cu @@ -20,7 +20,7 @@ #include #include #include -#include +#include #include #include #include @@ -124,13 +124,14 @@ namespace { */ struct integers_to_ipv4_fn { column_device_view const d_column; - size_type* d_offsets{}; + size_type* d_sizes{}; char* d_chars{}; + cudf::detail::input_offsetalator d_offsets; __device__ void operator()(size_type idx) { if (d_column.is_null(idx)) { - if (!d_chars) d_offsets[idx] = 0; + if (!d_chars) { d_sizes[idx] = 0; } return; } @@ -151,7 +152,7 @@ struct integers_to_ipv4_fn { shift_bits -= 8; } - if (!d_chars) { d_offsets[idx] = bytes; } + if (!d_chars) { d_sizes[idx] = bytes; } } }; @@ -167,7 +168,7 @@ std::unique_ptr integers_to_ipv4(column_view const& integers, CUDF_EXPECTS(integers.type().id() == type_id::INT64, "Input column must be type_id::INT64 type"); auto d_column = column_device_view::create(integers, stream); - auto [offsets_column, chars] = cudf::strings::detail::make_strings_children( + auto [offsets_column, chars] = experimental::make_strings_children( integers_to_ipv4_fn{*d_column}, integers.size(), stream, mr); return make_strings_column(integers.size(),