From 32c1f88339dd5e609d1f9e8cb47de6a939e204d1 Mon Sep 17 00:00:00 2001 From: Jason Lowe Date: Wed, 17 Jan 2024 11:22:06 -0600 Subject: [PATCH 1/2] Update to new cudf strings where character data is no longer a child column Signed-off-by: Jason Lowe --- src/main/cpp/src/cast_string_to_float.cu | 8 ++++---- src/main/cpp/src/map_utils.cu | 18 +++++++++++------- src/main/cpp/src/parse_uri.cu | 19 +++++++++++-------- thirdparty/cudf | 2 +- 4 files changed, 27 insertions(+), 20 deletions(-) diff --git a/src/main/cpp/src/cast_string_to_float.cu b/src/main/cpp/src/cast_string_to_float.cu index fe8a7f64db..75523cd360 100644 --- a/src/main/cpp/src/cast_string_to_float.cu +++ b/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 string_to_float(data_type dtype, out->mutable_view().null_mask(), ansi_mode ? static_cast(ansi_count.get())->data() : nullptr, static_cast(valid_count.get())->data(), - string_col.chars().begin(), + string_col.chars_begin(stream), string_col.offsets().begin(), string_col.null_mask(), num_rows); @@ -690,7 +690,7 @@ std::unique_ptr string_to_float(data_type dtype, out->mutable_view().null_mask(), ansi_mode ? static_cast(ansi_count.get())->data() : nullptr, static_cast(valid_count.get())->data(), - string_col.chars().begin(), + string_col.chars_begin(stream), string_col.offsets().begin(), string_col.null_mask(), num_rows); @@ -714,7 +714,7 @@ std::unique_ptr string_to_float(data_type dtype, dest.resize(string_bounds[1] - string_bounds[0]); cudaMemcpyAsync(dest.data(), - &string_col.chars().data()[string_bounds[0]], + &string_col.chars_begin(stream)[string_bounds[0]], string_bounds[1] - string_bounds[0], cudaMemcpyDeviceToHost, stream.value()); diff --git a/src/main/cpp/src/map_utils.cu b/src/main/cpp/src/map_utils.cu index a51a7de57b..761b0d94ea 100644 --- a/src/main/cpp/src/map_utils.cu +++ b/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,24 +75,28 @@ rmm::device_uvector 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(chars_size) + static_cast(input.size() - 1) + // append `,` character between input rows static_cast(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(std::numeric_limits::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 + "]". @@ -100,7 +104,7 @@ rmm::device_uvector unify_json_strings(cudf::column_view const& input, auto output = rmm::device_uvector(joined_input_size_bytes + 2, stream); CUDF_CUDA_TRY(cudaMemsetAsync(output.data(), static_cast('['), 1, stream.value())); CUDF_CUDA_TRY(cudaMemcpyAsync(output.data() + 1, - joined_input_child.view().data(), + joined_input_scv.chars_begin(stream), joined_input_size_bytes, cudaMemcpyDefault, stream.value())); diff --git a/src/main/cpp/src/parse_uri.cu b/src/main/cpp/src/parse_uri.cu index 897ebe0208..82ee044c2e 100644 --- a/src/main/cpp/src/parse_uri.cu +++ b/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(); 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(); for (thread_index_type tidx = tid; tidx < in_strings.size(); tidx += cudf::detail::grid_1d::grid_stride()) { @@ -840,6 +842,7 @@ std::unique_ptr parse_uri(strings_column_view const& input, parse_uri_char_counter<<>>( *d_strings, chunk, + input.chars_begin(stream), offsets_mutable_view.begin(), reinterpret_cast(src_offsets.data()), reinterpret_cast(null_mask.data())); @@ -854,23 +857,23 @@ std::unique_ptr parse_uri(strings_column_view const& input, // to the host memory auto out_chars_bytes = cudf::detail::get_value(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(); + // 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<<>>( *d_strings, + input.chars_begin(stream), reinterpret_cast(src_offsets.data()), offsets_column->view().begin(), - d_out_chars); + static_cast(d_out_chars.data())); auto null_count = cudf::null_count(reinterpret_cast(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)); } diff --git a/thirdparty/cudf b/thirdparty/cudf index 8f5e64ddcb..6abef4a474 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 8f5e64ddcba788ddcc715fda7f2bf852166b7ee6 +Subproject commit 6abef4a4746f1f9917711f372726023efdc21e85 From 38208e47559951ac4371aa29635acd01b0c975db Mon Sep 17 00:00:00 2001 From: Jason Lowe Date: Wed, 17 Jan 2024 11:29:14 -0600 Subject: [PATCH 2/2] clang style fixes --- src/main/cpp/src/map_utils.cu | 2 +- src/main/cpp/src/parse_uri.cu | 6 +++--- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/src/main/cpp/src/map_utils.cu b/src/main/cpp/src/map_utils.cu index 761b0d94ea..002dadb0e3 100644 --- a/src/main/cpp/src/map_utils.cu +++ b/src/main/cpp/src/map_utils.cu @@ -93,7 +93,7 @@ rmm::device_uvector unify_json_strings(cudf::column_view const& input, cudf::string_scalar("{}"), // replacement for null rows stream, rmm::mr::get_current_device_resource()); - auto const joined_input_scv = cudf::strings_column_view{*joined_input}; + 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 diff --git a/src/main/cpp/src/parse_uri.cu b/src/main/cpp/src/parse_uri.cu index 82ee044c2e..83b14ced9e 100644 --- a/src/main/cpp/src/parse_uri.cu +++ b/src/main/cpp/src/parse_uri.cu @@ -713,7 +713,7 @@ __global__ void parse_uri_char_counter(column_device_view const in_strings, bitmask_type* out_validity) { // thread per row - auto const tid = cudf::detail::grid_1d::global_thread_id(); + 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()) { @@ -790,7 +790,7 @@ __global__ void parse_uri(column_device_view const in_strings, size_type const* const offsets, char* const out_chars) { - auto const tid = cudf::detail::grid_1d::global_thread_id(); + 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()) { @@ -858,7 +858,7 @@ std::unique_ptr parse_uri(strings_column_view const& input, auto out_chars_bytes = cudf::detail::get_value(offsets_view, offset_count - 1, stream); // create the chars buffer - auto d_out_chars = rmm::device_buffer(out_chars_bytes, stream, mr); + 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<<>>(