From 507e33109b81df710cd3bdb2cfc71ddc35d759d6 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Mon, 17 Aug 2026 16:48:44 -0500 Subject: [PATCH 1/6] Migrate from rmm::cuda_stream_view to cuda::stream_ref Signed-off-by: Bradley Dice --- src/main/cpp/benchmarks/bloom_filter.cu | 6 +- .../benchmarks/cast_long_to_binary_string.cpp | 4 +- src/main/cpp/benchmarks/get_json_object.cu | 4 +- src/main/cpp/src/aggregation64_utils.cu | 4 +- src/main/cpp/src/aggregation64_utils.hpp | 8 +-- src/main/cpp/src/bloom_filter.cu | 22 +++---- src/main/cpp/src/bloom_filter.hpp | 13 ++-- src/main/cpp/src/case_when.cu | 4 +- src/main/cpp/src/case_when.hpp | 6 +- src/main/cpp/src/cast_decimal_to_string.cu | 10 +-- src/main/cpp/src/cast_float_to_string.cu | 13 ++-- .../cpp/src/cast_long_to_binary_string.cu | 12 ++-- src/main/cpp/src/cast_string.cu | 26 ++++---- src/main/cpp/src/cast_string.hpp | 26 ++++---- src/main/cpp/src/cast_string_to_datetime.cu | 10 +-- src/main/cpp/src/cast_string_to_float.cu | 10 +-- src/main/cpp/src/charset_decode.cu | 4 +- src/main/cpp/src/charset_decode.hpp | 5 +- src/main/cpp/src/datetime_rebase.cu | 12 ++-- src/main/cpp/src/datetime_truncate.cu | 14 ++-- src/main/cpp/src/datetime_utils.hpp | 12 ++-- src/main/cpp/src/decimal_utils.cu | 16 ++--- src/main/cpp/src/decimal_utils.hpp | 32 +++++----- .../src/exception_with_row_index_utilities.cu | 10 +-- .../exception_with_row_index_utilities.hpp | 14 ++-- src/main/cpp/src/format_float.cu | 10 +-- src/main/cpp/src/from_json_to_raw_map.cu | 50 +++++++-------- .../cpp/src/from_json_to_raw_map_debug.cuh | 13 ++-- src/main/cpp/src/from_json_to_structs.cu | 36 +++++------ src/main/cpp/src/get_json_object.cu | 16 ++--- src/main/cpp/src/get_json_object.hpp | 8 ++- src/main/cpp/src/hash/hash.hpp | 19 +++--- src/main/cpp/src/hash/hive_hash.cu | 4 +- src/main/cpp/src/hash/murmur_hash.cu | 4 +- src/main/cpp/src/hash/sha.cpp | 15 +++-- src/main/cpp/src/hash/xxhash64.cu | 4 +- src/main/cpp/src/hex.cu | 6 +- src/main/cpp/src/histogram.cu | 8 +-- src/main/cpp/src/histogram.hpp | 9 +-- src/main/cpp/src/hyper_log_log_plus_plus.hpp | 2 +- src/main/cpp/src/iceberg/iceberg_bucket.cu | 6 +- src/main/cpp/src/iceberg/iceberg_bucket.hpp | 6 +- .../cpp/src/iceberg/iceberg_datetime_util.cu | 16 ++--- .../cpp/src/iceberg/iceberg_datetime_util.hpp | 12 ++-- src/main/cpp/src/iceberg/iceberg_truncate.cu | 14 ++-- src/main/cpp/src/iceberg/iceberg_truncate.hpp | 10 +-- src/main/cpp/src/join_primitives.cu | 42 ++++++------ src/main/cpp/src/join_primitives.hpp | 21 +++--- src/main/cpp/src/json_utils.cu | 10 +-- src/main/cpp/src/json_utils.hpp | 15 +++-- src/main/cpp/src/list_slice.cu | 30 ++++----- src/main/cpp/src/list_slice.hpp | 12 ++-- src/main/cpp/src/map.cu | 4 +- src/main/cpp/src/map.hpp | 6 +- src/main/cpp/src/map_utils.cu | 58 ++++++++--------- src/main/cpp/src/map_utils.hpp | 6 +- src/main/cpp/src/map_zip_with_utils.cu | 7 +- src/main/cpp/src/map_zip_with_utils.hpp | 6 +- src/main/cpp/src/multiply.cu | 12 ++-- src/main/cpp/src/multiply.hpp | 10 +-- src/main/cpp/src/number_converter.cu | 10 +-- src/main/cpp/src/number_converter.hpp | 8 ++- .../cpp/src/parse_timestamp_with_format.cu | 6 +- src/main/cpp/src/parse_uri.cu | 22 +++---- src/main/cpp/src/parse_uri.hpp | 17 ++--- src/main/cpp/src/protobuf/protobuf.cu | 22 +++---- src/main/cpp/src/protobuf/protobuf.hpp | 5 +- .../cpp/src/protobuf/protobuf_builders.cu | 42 ++++++------ .../src/protobuf/protobuf_device_helpers.cuh | 5 +- .../src/protobuf/protobuf_host_helpers.hpp | 55 ++++++++-------- src/main/cpp/src/protobuf/protobuf_kernels.cu | 56 ++++++++-------- .../cpp/src/protobuf/protobuf_kernels.cuh | 45 ++++++------- src/main/cpp/src/regex_rewrite_utils.cu | 6 +- src/main/cpp/src/regex_rewrite_utils.hpp | 6 +- src/main/cpp/src/reverse_strings.cu | 6 +- src/main/cpp/src/reverse_strings.hpp | 4 +- src/main/cpp/src/round_float.cu | 14 ++-- src/main/cpp/src/round_float.hpp | 7 +- src/main/cpp/src/row_conversion.cu | 64 +++++++++---------- src/main/cpp/src/row_conversion.hpp | 13 ++-- src/main/cpp/src/shuffle_assemble.cu | 30 ++++----- src/main/cpp/src/shuffle_split.cu | 38 +++++------ src/main/cpp/src/shuffle_split.hpp | 6 +- src/main/cpp/src/substring_index.cu | 7 +- src/main/cpp/src/timezones.cu | 32 +++++----- src/main/cpp/src/timezones.hpp | 15 +++-- src/main/cpp/src/utilities.cu | 4 +- src/main/cpp/src/utilities.hpp | 7 +- src/main/cpp/src/uuid.cu | 6 +- src/main/cpp/src/uuid.hpp | 6 +- src/main/cpp/src/zorder.cu | 6 +- src/main/cpp/src/zorder.hpp | 9 +-- src/main/cpp/tests/hyper_log_log_plus_plus.cu | 9 +-- src/main/cpp/tests/protobuf_helpers.cu | 9 +-- src/main/cpp/tests/shuffle_split.cu | 2 +- 95 files changed, 723 insertions(+), 680 deletions(-) diff --git a/src/main/cpp/benchmarks/bloom_filter.cu b/src/main/cpp/benchmarks/bloom_filter.cu index 9f48c838d9..2ae8bc680d 100644 --- a/src/main/cpp/benchmarks/bloom_filter.cu +++ b/src/main/cpp/benchmarks/bloom_filter.cu @@ -18,6 +18,8 @@ #include +#include + #include #include #include @@ -40,12 +42,12 @@ void bloom_filter_put_impl(nvbench::state& state, int version) auto const input = spark_rapids_jni::xxhash64(*src); auto const stream = cudf::get_default_stream(); - state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value())); + state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.get())); state.exec(nvbench::exec_tag::timer | nvbench::exec_tag::sync, [&](nvbench::launch&, auto& timer) { timer.start(); spark_rapids_jni::bloom_filter_put(*bloom_filter, *input); - stream.synchronize(); + stream.sync(); timer.stop(); }); diff --git a/src/main/cpp/benchmarks/cast_long_to_binary_string.cpp b/src/main/cpp/benchmarks/cast_long_to_binary_string.cpp index c7179c0412..e95ed890d3 100644 --- a/src/main/cpp/benchmarks/cast_long_to_binary_string.cpp +++ b/src/main/cpp/benchmarks/cast_long_to_binary_string.cpp @@ -20,6 +20,8 @@ #include +#include + #include #include @@ -30,7 +32,7 @@ static void long_to_binary_string(nvbench::state& state) auto const input_table = create_random_table({cudf::type_id::INT64}, row_count{num_rows}); auto const long_col = input_table->get_column(0); auto const stream = cudf::get_default_stream(); - state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value())); + state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.get())); state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) { spark_rapids_jni::long_to_binary_string(long_col, stream); }); diff --git a/src/main/cpp/benchmarks/get_json_object.cu b/src/main/cpp/benchmarks/get_json_object.cu index b661c7c2e7..c78f6b321e 100644 --- a/src/main/cpp/benchmarks/get_json_object.cu +++ b/src/main/cpp/benchmarks/get_json_object.cu @@ -44,7 +44,7 @@ struct strings_to_host_fn { void operator()(std::vector& host_data, char const* chars, cudf::column_view const& offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const h_offsets = cudf::detail::make_std_vector_sync( cudf::device_span(offsets.data(), offsets.size()), stream); @@ -62,7 +62,7 @@ struct strings_to_host_fn { void operator()(std::vector&, char const*, cudf::column_view const&, - rmm::cuda_stream_view) + cuda::stream_ref) { CUDF_FAIL("invalid offsets type"); } diff --git a/src/main/cpp/src/aggregation64_utils.cu b/src/main/cpp/src/aggregation64_utils.cu index 3317185b99..93613f404b 100644 --- a/src/main/cpp/src/aggregation64_utils.cu +++ b/src/main/cpp/src/aggregation64_utils.cu @@ -75,7 +75,7 @@ namespace spark_rapids_jni { std::unique_ptr extract_chunk32_from_64bit(cudf::column_view const& in_col, cudf::data_type type, int chunk_idx, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS( @@ -112,7 +112,7 @@ std::unique_ptr extract_chunk32_from_64bit(cudf::column_view const // Reassemble a column of 64-bit values from two 64-bit integer columns with overflow detection. std::unique_ptr assemble64_from_sum(cudf::table_view const& chunks_table, cudf::data_type output_type, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS( diff --git a/src/main/cpp/src/aggregation64_utils.hpp b/src/main/cpp/src/aggregation64_utils.hpp index 66e0321a58..c504dd2233 100644 --- a/src/main/cpp/src/aggregation64_utils.hpp +++ b/src/main/cpp/src/aggregation64_utils.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2025, NVIDIA CORPORATION. + * Copyright (c) 2025-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -20,7 +20,7 @@ #include #include -#include +#include #include @@ -46,7 +46,7 @@ std::unique_ptr extract_chunk32_from_64bit( cudf::column_view const& col, cudf::data_type dtype, int chunk_idx, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -70,7 +70,7 @@ std::unique_ptr extract_chunk32_from_64bit( std::unique_ptr assemble64_from_sum( cudf::table_view const& chunks_table, cudf::data_type output_type, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/bloom_filter.cu b/src/main/cpp/src/bloom_filter.cu index f1271930ce..9be9987330 100644 --- a/src/main/cpp/src/bloom_filter.cu +++ b/src/main/cpp/src/bloom_filter.cu @@ -32,12 +32,12 @@ #include #include -#include #include #include #include #include +#include #include #include @@ -154,7 +154,7 @@ struct bloom_probe_functor { void pack_bloom_filter_header(cudf::device_span buf, bloom_filter_header const& header, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, int32_t seed) { if (header.version == bloom_filter_version_1) { @@ -188,7 +188,7 @@ void pack_bloom_filter_header(cudf::device_span buf, for V2, the value stored in the serialized header. */ std::tuple, int64_t, int32_t> -unpack_bloom_filter(cudf::device_span bloom_filter, rmm::cuda_stream_view stream) +unpack_bloom_filter(cudf::device_span bloom_filter, cuda::stream_ref stream) { CUDF_EXPECTS(bloom_filter.size() >= static_cast(bloom_filter_header_v1_size_bytes), "Encountered truncated bloom filter"); @@ -201,7 +201,7 @@ unpack_bloom_filter(cudf::device_span bloom_filter, rmm::cuda_str // Refer to https://github.com/NVIDIA/spark-rapids-jni/issues/4407. CUDF_CUDA_TRY( cudaMemcpyAsync(raw_ints, bloom_filter.data(), read_size, cudaMemcpyDefault, stream)); - stream.synchronize(); + stream.sync(); int const version = byte_swap_int32(raw_ints[0]); CUDF_EXPECTS(version == bloom_filter_version_1 || version == bloom_filter_version_2, @@ -237,7 +237,7 @@ unpack_bloom_filter(cudf::device_span bloom_filter, rmm::cuda_str } std::tuple, int64_t, int32_t> -unpack_bloom_filter(cudf::column_view const& bloom_filter, rmm::cuda_stream_view stream) +unpack_bloom_filter(cudf::column_view const& bloom_filter, cuda::stream_ref stream) { return unpack_bloom_filter( cudf::device_span{bloom_filter.data(), @@ -300,7 +300,7 @@ std::unique_ptr bloom_filter_create(int version, int num_hashes, int bloom_filter_longs, int seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -333,7 +333,7 @@ std::unique_ptr bloom_filter_create(int version, void bloom_filter_put(cudf::list_scalar& bloom_filter, cudf::column_view const& input, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { SRJ_FUNC_RANGE(); auto [header, buffer, bloom_filter_bits, seed] = unpack_bloom_filter(bloom_filter.view(), stream); @@ -351,7 +351,7 @@ void bloom_filter_put(cudf::list_scalar& bloom_filter, auto launch = [&](auto version_tag, auto nullable_tag) { gpu_bloom_filter_put - <<>>( + <<>>( mutable_buffer, bloom_filter_bits, *d_input, header.num_hashes, seed); }; @@ -373,7 +373,7 @@ void bloom_filter_put(cudf::list_scalar& bloom_filter, } std::unique_ptr bloom_filter_merge(cudf::column_view const& bloom_filters, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -450,7 +450,7 @@ std::unique_ptr bloom_filter_merge(cudf::column_view const& b std::unique_ptr bloom_filter_probe(cudf::column_view const& input, cudf::device_span bloom_filter, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -489,7 +489,7 @@ std::unique_ptr bloom_filter_probe(cudf::column_view const& input, std::unique_ptr bloom_filter_probe(cudf::column_view const& input, cudf::list_scalar& bloom_filter, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/bloom_filter.hpp b/src/main/cpp/src/bloom_filter.hpp index d6f413ca89..2bcae9a3dd 100644 --- a/src/main/cpp/src/bloom_filter.hpp +++ b/src/main/cpp/src/bloom_filter.hpp @@ -21,9 +21,10 @@ #include #include -#include #include +#include + namespace spark_rapids_jni { constexpr int bloom_filter_version_1 = 1; @@ -90,7 +91,7 @@ std::unique_ptr bloom_filter_create( int num_hashes, int bloom_filter_longs, int seed = 0, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -105,7 +106,7 @@ std::unique_ptr bloom_filter_create( */ void bloom_filter_put(cudf::list_scalar& bloom_filter, cudf::column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Probe a bloom filter with an input column of int64_t values. @@ -121,7 +122,7 @@ void bloom_filter_put(cudf::list_scalar& bloom_filter, std::unique_ptr bloom_filter_probe( cudf::column_view const& input, cudf::device_span bloom_filter, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -138,7 +139,7 @@ std::unique_ptr bloom_filter_probe( std::unique_ptr bloom_filter_probe( cudf::column_view const& input, cudf::list_scalar& bloom_filter, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -156,7 +157,7 @@ std::unique_ptr bloom_filter_probe( */ std::unique_ptr bloom_filter_merge( cudf::column_view const& bloom_filters, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/case_when.cu b/src/main/cpp/src/case_when.cu index 35abb6cfe4..a15cfb531d 100644 --- a/src/main/cpp/src/case_when.cu +++ b/src/main/cpp/src/case_when.cu @@ -67,7 +67,7 @@ struct select_first_true_fn { } // anonymous namespace std::unique_ptr select_first_true_index(cudf::table_view const& when_bool_columns, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // checks @@ -95,7 +95,7 @@ std::unique_ptr select_first_true_index(cudf::table_view const& wh } // namespace detail std::unique_ptr select_first_true_index(cudf::table_view const& when_bool_columns, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return detail::select_first_true_index(when_bool_columns, stream, mr); diff --git a/src/main/cpp/src/case_when.hpp b/src/main/cpp/src/case_when.hpp index b24b76608d..0f183fe3f7 100644 --- a/src/main/cpp/src/case_when.hpp +++ b/src/main/cpp/src/case_when.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2025, NVIDIA CORPORATION. + * Copyright (c) 2024-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -22,6 +22,8 @@ #include #include +#include + #include namespace spark_rapids_jni { @@ -47,7 +49,7 @@ namespace spark_rapids_jni { */ std::unique_ptr select_first_true_index( cudf::table_view const& when_bool_columns, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/cast_decimal_to_string.cu b/src/main/cpp/src/cast_decimal_to_string.cu index a80339a3a8..4b89f38cca 100644 --- a/src/main/cpp/src/cast_decimal_to_string.cu +++ b/src/main/cpp/src/cast_decimal_to_string.cu @@ -29,13 +29,13 @@ #include #include -#include #include #include #include #include #include +#include #include #include #include @@ -180,7 +180,7 @@ struct decimal_to_non_ansi_string_fn { struct dispatch_decimal_to_non_ansi_string_fn { template ()>* = nullptr> std::unique_ptr operator()(column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { using DecimalType = device_storage_type_t; // underlying value type @@ -199,7 +199,7 @@ struct dispatch_decimal_to_non_ansi_string_fn { template ()>* = nullptr> std::unique_ptr operator()(column_view const&, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref) const { CUDF_FAIL("Values for decimal_to_non_ansi_string function must be a decimal type."); @@ -209,7 +209,7 @@ struct dispatch_decimal_to_non_ansi_string_fn { } // namespace std::unique_ptr decimal_to_non_ansi_string(column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.is_empty()) return make_empty_column(type_id::STRING); @@ -221,7 +221,7 @@ std::unique_ptr decimal_to_non_ansi_string(column_view const& input, // external API std::unique_ptr decimal_to_non_ansi_string(column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/cast_float_to_string.cu b/src/main/cpp/src/cast_float_to_string.cu index cde21575f5..d388cc065a 100644 --- a/src/main/cpp/src/cast_float_to_string.cu +++ b/src/main/cpp/src/cast_float_to_string.cu @@ -24,9 +24,10 @@ #include #include -#include #include +#include + namespace spark_rapids_jni { namespace detail { @@ -76,7 +77,7 @@ template struct dispatch_float_to_string_fn { template )> std::unique_ptr operator()(cudf::column_view const& floats, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const strings_count = floats.size(); @@ -98,7 +99,7 @@ struct dispatch_float_to_string_fn { // non-float types throw an exception template )> std::unique_ptr operator()(cudf::column_view const&, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref) { CUDF_FAIL("Values for float_to_string function must be a float type."); @@ -110,7 +111,7 @@ struct dispatch_float_to_string_fn { // This will convert all float column types into a strings column. std::unique_ptr float_to_string(cudf::column_view const& floats, bool json_string, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return json_string @@ -123,7 +124,7 @@ std::unique_ptr float_to_string(cudf::column_view const& floats, // external API std::unique_ptr float_to_string(cudf::column_view const& floats, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return float_to_string(floats, false, stream, mr); @@ -131,7 +132,7 @@ std::unique_ptr float_to_string(cudf::column_view const& floats, std::unique_ptr float_to_string(cudf::column_view const& floats, bool json_string, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/cast_long_to_binary_string.cu b/src/main/cpp/src/cast_long_to_binary_string.cu index 13dcfe737a..5754a61be5 100644 --- a/src/main/cpp/src/cast_long_to_binary_string.cu +++ b/src/main/cpp/src/cast_long_to_binary_string.cu @@ -24,9 +24,10 @@ #include #include -#include #include +#include + namespace spark_rapids_jni { namespace detail { @@ -96,7 +97,7 @@ CUDF_KERNEL void long_to_binary_string_kernel(cudf::column_device_view d_longs, } // namespace std::unique_ptr long_to_binary_string(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input.type().id() == cudf::type_id::INT64, "Input column must be long type"); @@ -113,8 +114,7 @@ std::unique_ptr long_to_binary_string(cudf::column_view const& inp cudf::size_type* d_sizes = output_sizes.data(); auto constexpr block_size = 256; auto grid = cudf::detail::grid_1d{strings_count, block_size}; - compute_output_size_kernel<<>>(*d_column, - d_sizes); + compute_output_size_kernel<<>>(*d_column, d_sizes); // Convert the sizes to offsets auto [offsets, bytes] = cudf::strings::detail::make_offsets_child_column( @@ -131,7 +131,7 @@ std::unique_ptr long_to_binary_string(cudf::column_view const& inp auto new_grid = cudf::detail::grid_1d{strings_count * num_threads_per_row, block_size}; if (bytes > 0) { long_to_binary_string_kernel - <<>>(*d_column, d_chars, d_offsets); + <<>>(*d_column, d_chars, d_offsets); } return cudf::make_strings_column(input.size(), @@ -145,7 +145,7 @@ std::unique_ptr long_to_binary_string(cudf::column_view const& inp // external API std::unique_ptr long_to_binary_string(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/cast_string.cu b/src/main/cpp/src/cast_string.cu index 046b60f94b..1f35235ec4 100644 --- a/src/main/cpp/src/cast_string.cu +++ b/src/main/cpp/src/cast_string.cu @@ -610,7 +610,7 @@ struct row_valid_fn { */ void validate_ansi_column(column_view const& col, strings_column_view const& source_col, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const num_nulls = col.null_count(); auto const incoming_nulls = source_col.null_count(); @@ -627,8 +627,8 @@ void validate_ansi_column(column_view const& col, &source_col.offsets().data()[*first_error], sizeof(size_type) * 2, cudaMemcpyDefault, - stream.value()); - stream.synchronize(); + stream.get()); + stream.sync(); std::string dest; dest.resize(string_bounds[1] - string_bounds[0]); @@ -637,8 +637,8 @@ void validate_ansi_column(column_view const& col, &source_col.chars_begin(stream)[string_bounds[0]], string_bounds[1] - string_bounds[0], cudaMemcpyDefault, - stream.value()); - stream.synchronize(); + stream.get()); + stream.sync(); throw cast_error(*first_error, dest); } @@ -659,7 +659,7 @@ struct string_to_integer_impl { std::unique_ptr operator()(strings_column_view const& string_col, bool ansi_mode, bool strip, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (string_col.size() == 0) { @@ -674,7 +674,7 @@ struct string_to_integer_impl { dim3 const blocks(util::div_rounding_up_unsafe(string_col.size(), detail::NUM_THREADS)); dim3 const threads{detail::NUM_THREADS}; - detail::string_to_integer_kernel<<>>( + detail::string_to_integer_kernel<<>>( data.data(), null_mask.data(), string_col.chars_begin(stream), @@ -704,7 +704,7 @@ struct string_to_integer_impl { std::unique_ptr operator()(strings_column_view const& string_col, bool ansi_mode, bool strip, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FAIL("Invalid integer column type"); @@ -731,7 +731,7 @@ struct string_to_decimal_impl { strings_column_view const& string_col, bool ansi_mode, bool strip, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { using Type = device_storage_type_t; @@ -743,7 +743,7 @@ struct string_to_decimal_impl { dim3 const blocks(util::div_rounding_up_unsafe(string_col.size(), detail::NUM_THREADS)); dim3 const threads{detail::NUM_THREADS}; - detail::string_to_decimal_kernel<<>>( + detail::string_to_decimal_kernel<<>>( data.data(), null_mask.data(), string_col.chars_begin(stream), @@ -773,7 +773,7 @@ struct string_to_decimal_impl { strings_column_view const& string_col, bool ansi_mode, bool strip, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FAIL("Invalid decimal column type"); @@ -798,7 +798,7 @@ std::unique_ptr string_to_integer(data_type dtype, strings_column_view const& string_col, bool ansi_mode, bool strip, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return type_dispatcher( @@ -823,7 +823,7 @@ std::unique_ptr string_to_decimal(int32_t precision, strings_column_view const& string_col, bool ansi_mode, bool strip, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { data_type dtype = [precision, scale]() { diff --git a/src/main/cpp/src/cast_string.hpp b/src/main/cpp/src/cast_string.hpp index 9f4ce4cd6f..767e4633a3 100644 --- a/src/main/cpp/src/cast_string.hpp +++ b/src/main/cpp/src/cast_string.hpp @@ -25,6 +25,8 @@ #include +#include + #include namespace spark_rapids_jni { @@ -78,7 +80,7 @@ std::unique_ptr string_to_integer( cudf::strings_column_view const& string_col, bool ansi_mode, bool strip, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -100,7 +102,7 @@ std::unique_ptr string_to_decimal( cudf::strings_column_view const& string_col, bool ansi_mode, bool strip, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -118,34 +120,34 @@ std::unique_ptr string_to_float( cudf::data_type dtype, cudf::strings_column_view const& string_col, bool ansi_mode, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); std::unique_ptr format_float( cudf::column_view const& input, int const digits, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); std::unique_ptr float_to_string( cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); [[nodiscard]] std::unique_ptr float_to_string( cudf::column_view const& input, bool json_string, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); std::unique_ptr decimal_to_non_ansi_string( cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); std::unique_ptr long_to_binary_string( cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -176,7 +178,7 @@ std::unique_ptr parse_timestamp_strings( cudf::column_view const& tz_name_to_index_map, cudf::table_view const& tz_info_table, spark_rapids_jni::spark_system const& spark_system, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -198,7 +200,7 @@ std::unique_ptr parse_timestamp_strings( */ std::unique_ptr parse_strings_to_date( cudf::strings_column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -237,7 +239,7 @@ std::unique_ptr parse_timestamp_strings_with_format( cudf::strings_column_view const& input, std::string const& format, bool legacy, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -253,7 +255,7 @@ std::unique_ptr parse_timestamp_strings_with_format( */ std::unique_ptr bytes_to_hex( cudf::strings_column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/cast_string_to_datetime.cu b/src/main/cpp/src/cast_string_to_datetime.cu index fad6a96067..65ef54aa20 100644 --- a/src/main/cpp/src/cast_string_to_datetime.cu +++ b/src/main/cpp/src/cast_string_to_datetime.cu @@ -27,10 +27,10 @@ #include #include -#include #include #include +#include #include #include #include @@ -875,7 +875,7 @@ std::unique_ptr parse_ts_strings(cudf::strings_column_view const& cudf::table_view const& tz_info_table, bool is_spark_320, bool is_spark_400_or_later_or_db_14_3_or_later, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const num_rows = input.size(); @@ -1078,7 +1078,7 @@ struct parse_string_to_date_fn { * Parse strings to dates. */ std::unique_ptr parse_to_date(cudf::strings_column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const num_rows = input.size(); @@ -1120,7 +1120,7 @@ std::unique_ptr parse_timestamp_strings( cudf::column_view const& tz_name_to_index_map, cudf::table_view const& tz_info_table, spark_rapids_jni::spark_system const& spark_system, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { bool is_spark_320 = spark_system.is_vanilla_320(); @@ -1139,7 +1139,7 @@ std::unique_ptr parse_timestamp_strings( } std::unique_ptr parse_strings_to_date(cudf::strings_column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return parse_to_date(input, stream, mr); diff --git a/src/main/cpp/src/cast_string_to_float.cu b/src/main/cpp/src/cast_string_to_float.cu index 8515d8c238..42461bbd3a 100644 --- a/src/main/cpp/src/cast_string_to_float.cu +++ b/src/main/cpp/src/cast_string_to_float.cu @@ -886,7 +886,7 @@ CUDF_KERNEL void string_to_float_kernel(T* out, std::unique_ptr string_to_float(data_type dtype, strings_column_view const& string_col, bool ansi_mode, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(dtype == data_type{type_id::FLOAT32} || dtype == data_type{type_id::FLOAT64}, @@ -941,8 +941,8 @@ std::unique_ptr string_to_float(data_type dtype, &string_col.offsets().data()[error_row], sizeof(size_type) * 2, cudaMemcpyDefault, - stream.value()); - stream.synchronize(); + stream.get()); + stream.sync(); std::string dest; dest.resize(string_bounds[1] - string_bounds[0]); @@ -951,8 +951,8 @@ std::unique_ptr string_to_float(data_type dtype, &string_col.chars_begin(stream)[string_bounds[0]], string_bounds[1] - string_bounds[0], cudaMemcpyDefault, - stream.value()); - stream.synchronize(); + stream.get()); + stream.sync(); throw cast_error(error_row, dest); } diff --git a/src/main/cpp/src/charset_decode.cu b/src/main/cpp/src/charset_decode.cu index 7ecd007676..dfda595645 100644 --- a/src/main/cpp/src/charset_decode.cu +++ b/src/main/cpp/src/charset_decode.cu @@ -202,7 +202,7 @@ struct gbk_decode_fn { decode_result decode_gbk(cudf::column_view const& input, error_action action, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input.type().id() == cudf::type_id::LIST, @@ -256,7 +256,7 @@ decode_result decode_gbk(cudf::column_view const& input, decode_result decode_charset(cudf::column_view const& input, charset_type charset, error_action action, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/charset_decode.hpp b/src/main/cpp/src/charset_decode.hpp index 47b239d244..dc769ca34f 100644 --- a/src/main/cpp/src/charset_decode.hpp +++ b/src/main/cpp/src/charset_decode.hpp @@ -21,9 +21,10 @@ #include #include -#include #include +#include + #include namespace spark_rapids_jni { @@ -77,7 +78,7 @@ struct decode_result { cudf::column_view const& input, charset_type charset, error_action action, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/datetime_rebase.cu b/src/main/cpp/src/datetime_rebase.cu index 4ced40b710..3955f5c078 100644 --- a/src/main/cpp/src/datetime_rebase.cu +++ b/src/main/cpp/src/datetime_rebase.cu @@ -53,7 +53,7 @@ __device__ __inline__ auto days_from_julian(cuda::std::chrono::year_month_day co // days since the epoch from that Julian local date. // This is to match with Apache Spark's `localRebaseGregorianToJulianDays` function. std::unique_ptr gregorian_to_julian_days(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input.type().id() == cudf::type_id::TIMESTAMP_DAYS, @@ -124,7 +124,7 @@ __device__ __inline__ cuda::std::chrono::year_month_day julian_from_days(int32_t // of days since the epoch from that Gregorian local date. This is to match with Apache Spark's // `localRebaseJulianToGregorianDays` function. std::unique_ptr julian_to_gregorian_days(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input.type().id() == cudf::type_id::TIMESTAMP_DAYS, @@ -224,7 +224,7 @@ __device__ __inline__ time_components get_time_components(int64_t micros) // This is to match with Apache Spark's `rebaseGregorianToJulianMicros` function with timezone // fixed to UTC. std::unique_ptr gregorian_to_julian_micros(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input.type().id() == cudf::type_id::TIMESTAMP_MICROSECONDS, @@ -287,7 +287,7 @@ std::unique_ptr gregorian_to_julian_micros(cudf::column_view const // This is to match with Apache Spark's `rebaseJulianToGregorianMicros` function with timezone // fixed to UTC. std::unique_ptr julian_to_gregorian_micros(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input.type().id() == cudf::type_id::TIMESTAMP_MICROSECONDS, @@ -341,7 +341,7 @@ std::unique_ptr julian_to_gregorian_micros(cudf::column_view const namespace spark_rapids_jni { std::unique_ptr rebase_gregorian_to_julian(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -357,7 +357,7 @@ std::unique_ptr rebase_gregorian_to_julian(cudf::column_view const } std::unique_ptr rebase_julian_to_gregorian(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/datetime_truncate.cu b/src/main/cpp/src/datetime_truncate.cu index 363d44aaa8..e024f5082b 100644 --- a/src/main/cpp/src/datetime_truncate.cu +++ b/src/main/cpp/src/datetime_truncate.cu @@ -26,13 +26,13 @@ #include #include -#include #include #include #include #include #include +#include #include #include #include @@ -266,7 +266,7 @@ template std::unique_ptr truncate_datetime(cudf::column_view const& datetime, FormatT const& format, cudf::size_type output_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto output = cudf::make_fixed_width_column( @@ -318,7 +318,7 @@ template std::unique_ptr truncate_dispatcher(cudf::column_view const& datetime, FormatT const& format, cudf::size_type output_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (datetime.type().id() == cudf::type_id::TIMESTAMP_DAYS) { @@ -345,7 +345,7 @@ void check_types(cudf::column_view const& datetime, cudf::column_view const& for std::unique_ptr truncate(cudf::column_view const& datetime, cudf::column_view const& format, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { check_types(datetime, format); @@ -364,7 +364,7 @@ std::unique_ptr truncate(cudf::column_view const& datetime, std::unique_ptr truncate(cudf::column_view const& datetime, std::string const& format, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { check_type(datetime); @@ -382,7 +382,7 @@ std::unique_ptr truncate(cudf::column_view const& datetime, std::unique_ptr truncate(cudf::column_view const& datetime, cudf::column_view const& format, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -391,7 +391,7 @@ std::unique_ptr truncate(cudf::column_view const& datetime, std::unique_ptr truncate(cudf::column_view const& datetime, std::string const& format, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/datetime_utils.hpp b/src/main/cpp/src/datetime_utils.hpp index b531317475..144b8fd6df 100644 --- a/src/main/cpp/src/datetime_utils.hpp +++ b/src/main/cpp/src/datetime_utils.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * Copyright (c) 2023-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -18,27 +18,29 @@ #include +#include + namespace spark_rapids_jni { std::unique_ptr rebase_gregorian_to_julian( cudf::column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); std::unique_ptr rebase_julian_to_gregorian( cudf::column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); std::unique_ptr truncate( cudf::column_view const& datetime, cudf::column_view const& format, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); std::unique_ptr truncate( cudf::column_view const& datetime, std::string const& format, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/decimal_utils.cu b/src/main/cpp/src/decimal_utils.cu index 877f9817b7..7af5518b54 100644 --- a/src/main/cpp/src/decimal_utils.cu +++ b/src/main/cpp/src/decimal_utils.cu @@ -969,7 +969,7 @@ std::unique_ptr multiply_decimal128(cudf::column_view const& a, cudf::column_view const& b, int32_t product_scale, bool const cast_interim_result, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS(a.type().id() == cudf::type_id::DECIMAL128, "not a DECIMAL128 column"); CUDF_EXPECTS(b.type().id() == cudf::type_id::DECIMAL128, "not a DECIMAL128 column"); @@ -1004,7 +1004,7 @@ std::unique_ptr multiply_decimal128(cudf::column_view const& a, std::unique_ptr divide_decimal128(cudf::column_view const& a, cudf::column_view const& b, int32_t quotient_scale, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS(a.type().id() == cudf::type_id::DECIMAL128, "not a DECIMAL128 column"); CUDF_EXPECTS(b.type().id() == cudf::type_id::DECIMAL128, "not a DECIMAL128 column"); @@ -1038,7 +1038,7 @@ std::unique_ptr divide_decimal128(cudf::column_view const& a, std::unique_ptr integer_divide_decimal128(cudf::column_view const& a, cudf::column_view const& b, int32_t quotient_scale, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS(a.type().id() == cudf::type_id::DECIMAL128, "not a DECIMAL128 column"); CUDF_EXPECTS(b.type().id() == cudf::type_id::DECIMAL128, "not a DECIMAL128 column"); @@ -1071,7 +1071,7 @@ std::unique_ptr integer_divide_decimal128(cudf::column_view const& std::unique_ptr remainder_decimal128(cudf::column_view const& a, cudf::column_view const& b, int32_t remainder_scale, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS(a.type().id() == cudf::type_id::DECIMAL128, "not a DECIMAL128 column"); CUDF_EXPECTS(b.type().id() == cudf::type_id::DECIMAL128, "not a DECIMAL128 column"); @@ -1104,7 +1104,7 @@ std::unique_ptr remainder_decimal128(cudf::column_view const& a, std::unique_ptr add_decimal128(cudf::column_view const& a, cudf::column_view const& b, int32_t target_scale, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS(a.type().id() == cudf::type_id::DECIMAL128, "not a DECIMAL128 column"); CUDF_EXPECTS(b.type().id() == cudf::type_id::DECIMAL128, "not a DECIMAL128 column"); @@ -1137,7 +1137,7 @@ std::unique_ptr add_decimal128(cudf::column_view const& a, std::unique_ptr sub_decimal128(cudf::column_view const& a, cudf::column_view const& b, int32_t target_scale, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS(a.type().id() == cudf::type_id::DECIMAL128, "not a DECIMAL128 column"); CUDF_EXPECTS(b.type().id() == cudf::type_id::DECIMAL128, "not a DECIMAL128 column"); @@ -1365,7 +1365,7 @@ struct floating_point_to_decimal_dispatcher { cudf::size_type* failure_row_id, int32_t decimal_places, int32_t precision, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { using DecimalRepType = cudf::device_storage_type_t; @@ -1388,7 +1388,7 @@ std::pair, cudf::size_type> floating_point_to_deci cudf::column_view const& input, cudf::data_type output_type, int32_t precision, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto output = cudf::make_fixed_point_column( diff --git a/src/main/cpp/src/decimal_utils.hpp b/src/main/cpp/src/decimal_utils.hpp index 91295cfdfa..562bd5046a 100644 --- a/src/main/cpp/src/decimal_utils.hpp +++ b/src/main/cpp/src/decimal_utils.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2025, NVIDIA CORPORATION. + * Copyright (c) 2022-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -20,7 +20,7 @@ #include #include -#include +#include #include @@ -31,37 +31,35 @@ std::unique_ptr multiply_decimal128( cudf::column_view const& b, int32_t product_scale, bool const cast_interim_result, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); std::unique_ptr divide_decimal128( cudf::column_view const& a, cudf::column_view const& b, int32_t quotient_scale, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); std::unique_ptr integer_divide_decimal128( cudf::column_view const& a, cudf::column_view const& b, int32_t quotient_scale, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); std::unique_ptr remainder_decimal128( cudf::column_view const& a, cudf::column_view const& b, int32_t remainder_scale, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); -std::unique_ptr add_decimal128( - cudf::column_view const& a, - cudf::column_view const& b, - int32_t quotient_scale, - rmm::cuda_stream_view stream = cudf::get_default_stream()); +std::unique_ptr add_decimal128(cudf::column_view const& a, + cudf::column_view const& b, + int32_t quotient_scale, + cuda::stream_ref stream = cudf::get_default_stream()); -std::unique_ptr sub_decimal128( - cudf::column_view const& a, - cudf::column_view const& b, - int32_t quotient_scale, - rmm::cuda_stream_view stream = cudf::get_default_stream()); +std::unique_ptr sub_decimal128(cudf::column_view const& a, + cudf::column_view const& b, + int32_t quotient_scale, + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Cast floating point values to decimals, matching the behavior of Spark. @@ -78,7 +76,7 @@ std::pair, cudf::size_type> floating_point_to_deci cudf::column_view const& input, cudf::data_type output_type, int32_t precision, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); } // namespace cudf::jni diff --git a/src/main/cpp/src/exception_with_row_index_utilities.cu b/src/main/cpp/src/exception_with_row_index_utilities.cu index a39ece0771..eaac777bc7 100644 --- a/src/main/cpp/src/exception_with_row_index_utilities.cu +++ b/src/main/cpp/src/exception_with_row_index_utilities.cu @@ -91,7 +91,7 @@ struct row_invalid_ternary_fn { void throw_row_error_if_any(cudf::column_view const& input, cudf::column_view const& result, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS(input.size() == result.size(), "The row counts of the input and result columns must match."); @@ -115,7 +115,7 @@ void throw_row_error_if_any(cudf::column_view const& input, void throw_row_error_if_any(cudf::column_view const& input1, cudf::column_view const& input2, cudf::column_view const& result, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS((input1.size() == input2.size() && input2.size() == result.size()), "The row counts of the input and result columns must match."); @@ -138,7 +138,7 @@ void throw_row_error_if_any(cudf::column_view const& input1, cudf::column_view const& input2, cudf::column_view const& input3, cudf::column_view const& result, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS((input1.size() == input2.size() && input2.size() == input3.size() && input3.size() == result.size()), @@ -162,7 +162,7 @@ void throw_row_error_if_any(cudf::column_view const& input1, void throw_row_error_if_any(cudf::column_view const& input1, cudf::scalar const& input2, cudf::column_view const& result, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS(input1.size() == result.size(), "The row counts of the input and result columns must match."); @@ -180,7 +180,7 @@ void throw_row_error_if_any(cudf::column_view const& input1, void throw_row_error_if_any(cudf::scalar const& input1, cudf::column_view const& input2, cudf::column_view const& result, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { throw_row_error_if_any(input2, input1, result, stream); } diff --git a/src/main/cpp/src/exception_with_row_index_utilities.hpp b/src/main/cpp/src/exception_with_row_index_utilities.hpp index e35e81ad25..39dbfb87fe 100644 --- a/src/main/cpp/src/exception_with_row_index_utilities.hpp +++ b/src/main/cpp/src/exception_with_row_index_utilities.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2025, NVIDIA CORPORATION. + * Copyright (c) 2025-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -19,6 +19,8 @@ #include #include +#include + namespace spark_rapids_jni { /** @@ -32,7 +34,7 @@ namespace spark_rapids_jni { */ void throw_row_error_if_any(cudf::column_view const& input, cudf::column_view const& result, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Throws exception_with_row_index if has any row is invalid for a binary operation. @@ -47,7 +49,7 @@ void throw_row_error_if_any(cudf::column_view const& input, void throw_row_error_if_any(cudf::column_view const& input1, cudf::column_view const& input2, cudf::column_view const& result, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Throws exception_with_row_index if has any row is invalid for a ternary operation. @@ -64,7 +66,7 @@ void throw_row_error_if_any(cudf::column_view const& input1, cudf::column_view const& input2, cudf::column_view const& input3, cudf::column_view const& result, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Throws exception_with_row_index if has any row is invalid for a binary operation. @@ -79,7 +81,7 @@ void throw_row_error_if_any(cudf::column_view const& input1, void throw_row_error_if_any(cudf::column_view const& input1, cudf::scalar const& input2, cudf::column_view const& result, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Throws exception_with_row_index if has any row is invalid for a binary operation. @@ -94,6 +96,6 @@ void throw_row_error_if_any(cudf::column_view const& input1, void throw_row_error_if_any(cudf::scalar const& input1, cudf::column_view const& input2, cudf::column_view const& result, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/format_float.cu b/src/main/cpp/src/format_float.cu index ae0fb9a540..7a5ee5449d 100644 --- a/src/main/cpp/src/format_float.cu +++ b/src/main/cpp/src/format_float.cu @@ -24,10 +24,10 @@ #include #include -#include #include #include +#include namespace spark_rapids_jni { @@ -80,7 +80,7 @@ struct dispatch_format_float_fn { template )> std::unique_ptr operator()(cudf::column_view const& floats, int const digits, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { auto const strings_count = floats.size(); @@ -103,7 +103,7 @@ struct dispatch_format_float_fn { template )> std::unique_ptr operator()(cudf::column_view const&, int const, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref) const { CUDF_FAIL("Values for format_float function must be a float type."); @@ -115,7 +115,7 @@ struct dispatch_format_float_fn { // This will convert all float column types into a strings column. std::unique_ptr format_float(cudf::column_view const& floats, int const digits, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return type_dispatcher(floats.type(), dispatch_format_float_fn{}, floats, digits, stream, mr); @@ -126,7 +126,7 @@ std::unique_ptr format_float(cudf::column_view const& floats, // external API std::unique_ptr format_float(cudf::column_view const& floats, int const digits, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/from_json_to_raw_map.cu b/src/main/cpp/src/from_json_to_raw_map.cu index 1f3a9f6c5e..72f8f5534b 100644 --- a/src/main/cpp/src/from_json_to_raw_map.cu +++ b/src/main/cpp/src/from_json_to_raw_map.cu @@ -30,7 +30,6 @@ #include #include -#include #include #include #include @@ -42,6 +41,7 @@ #include #include #include +#include #include #include #include @@ -69,7 +69,7 @@ OutputIterator copy_if(InputIterator begin, StencilIterator stencil, OutputIterator result, Predicate predicate, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const num_items = cuda::std::distance(begin, end); @@ -85,7 +85,7 @@ OutputIterator copy_if(InputIterator begin, num_selected.data(), num_items, predicate, - stream.value())); + stream.get())); auto d_temp_storage = rmm::device_buffer(temp_storage_bytes, stream, cudf::get_current_device_resource_ref()); @@ -98,7 +98,7 @@ OutputIterator copy_if(InputIterator begin, num_selected.data(), num_items, predicate, - stream.value())); + stream.get())); return result + num_selected.value(stream); } @@ -108,7 +108,7 @@ OutputIterator copy_if(InputIterator begin, InputIterator end, OutputIterator output, Predicate predicate, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const num_items = cuda::std::distance(begin, end); @@ -125,7 +125,7 @@ OutputIterator copy_if(InputIterator begin, num_selected.data(), num_items, predicate, - stream.value())); + stream.get())); // Allocate temporary storage rmm::device_buffer d_temp_storage( @@ -139,7 +139,7 @@ OutputIterator copy_if(InputIterator begin, num_selected.data(), num_items, predicate, - stream.value())); + stream.get())); // Copy number of selected elements back to host via pinned memory return output + num_selected.value(stream); @@ -149,7 +149,7 @@ OutputIterator copy_if(InputIterator begin, // value-child type selects the map flavor: an empty `STRING` gives `make_empty_map`'s output, an // empty `List` gives `make_empty_map_array`'s. std::unique_ptr make_empty_map_from_value(std::unique_ptr value_child, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(value_child->size() == 0, "value_child must be an empty column."); @@ -163,7 +163,7 @@ std::unique_ptr make_empty_map_from_value(std::unique_ptr make_empty_map(rmm::cuda_stream_view stream, +std::unique_ptr make_empty_map(cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return make_empty_map_from_value( @@ -173,7 +173,7 @@ std::unique_ptr make_empty_map(rmm::cuda_stream_view stream, // Concatenating all input strings into one string, for which each input string is appended by a // delimiter character that does not exist in the input column. std::tuple> unify_json_strings( - cudf::strings_column_view const& input, rmm::cuda_stream_view stream) + cudf::strings_column_view const& input, cuda::stream_ref stream) { auto const default_mr = cudf::get_current_device_resource_ref(); auto [concatenated_buff, delimiter, should_be_nullified] = @@ -224,7 +224,7 @@ struct is_node { // This is copied from cudf's `json_tree.cu`. rmm::device_uvector compute_node_levels(std::size_t num_nodes, cudf::device_span tokens, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto token_levels = rmm::device_uvector(tokens.size(), stream); @@ -279,7 +279,7 @@ rmm::device_uvector compute_node_levels(std::size_t num_nodes, // Compute the map from nodes to their indices in the list of all tokens. rmm::device_uvector compute_node_to_token_index_map( - std::size_t num_nodes, cudf::device_span tokens, rmm::cuda_stream_view stream) + std::size_t num_nodes, cudf::device_span tokens, cuda::stream_ref stream) { auto node_token_ids = rmm::device_uvector(num_nodes, stream); auto const node_id_it = thrust::counting_iterator(0); @@ -301,7 +301,7 @@ rmm::device_uvector compute_node_to_token_index_map( // This is copied from cudf's `json_tree.cu`. template std::pair, rmm::device_uvector> stable_sorted_key_order( - cudf::device_span keys, rmm::cuda_stream_view stream) + cudf::device_span keys, cuda::stream_ref stream) { // Buffers used for storing intermediate results during sorting. rmm::device_uvector keys_buffer1(keys.size(), stream); @@ -330,7 +330,7 @@ std::pair, rmm::device_uvector> stable_s keys.size(), 0, sizeof(KeyType) * 8, - stream.value()); + stream.get()); return std::pair{keys_buffer.Current() == keys_buffer1.data() ? std::move(keys_buffer1) : std::move(keys_buffer2), @@ -341,7 +341,7 @@ std::pair, rmm::device_uvector> stable_s // This is copied from cudf's `json_tree.cu`. void propagate_parent_to_siblings(cudf::device_span node_levels, cudf::device_span parent_node_ids, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const [sorted_node_levels, sorted_order] = stable_sorted_key_order(node_levels, stream); @@ -360,7 +360,7 @@ void propagate_parent_to_siblings(cudf::device_span node_level rmm::device_uvector compute_parent_node_ids( cudf::device_span tokens, cudf::device_span node_token_ids, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const first_childs_parent_token_id = cuda::proclaim_return_type([tokens] __device__(auto i) -> NodeIndexT { @@ -423,7 +423,7 @@ constexpr int32_t list_nesting_weight{1 << 8}; // Check for each node if it is a key or a value field. rmm::device_uvector check_key_or_value_nodes( - cudf::device_span parent_node_ids, rmm::cuda_stream_view stream) + cudf::device_span parent_node_ids, cuda::stream_ref stream) { auto key_or_value = rmm::device_uvector(parent_node_ids.size(), stream); auto const transform_it = thrust::counting_iterator(0); @@ -475,7 +475,7 @@ struct tokenized_input { // this and diverge only after it returns. tokenized_input tokenize_and_classify(cudf::strings_column_view const& input, json_parse_options const& options, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto [concat_json_buff, delimiter, should_be_nullified] = unify_json_strings(input, stream); auto concat_buff_wrapper = @@ -675,7 +675,7 @@ rmm::device_uvector> compute_node_ cudf::device_span node_token_ids, cudf::device_span parent_node_ids, cudf::device_span key_or_value, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const num_nodes = node_token_ids.size(); auto node_ranges = @@ -723,7 +723,7 @@ std::unique_ptr extract_keys_or_values( cudf::device_span const> node_ranges, cudf::device_span key_or_value, cudf::device_span input_json, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const is_key_or_value = cuda::proclaim_return_type( @@ -753,7 +753,7 @@ std::unique_ptr compute_list_offsets( cudf::size_type n_lists, cudf::device_span parent_node_ids, cudf::device_span key_or_value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // Count the number of children nodes for the json object nodes. @@ -859,7 +859,7 @@ std::pair create_null_mask( cudf::device_span node_token_ids, cudf::device_span parent_node_ids, cudf::device_span precomputed_line_begin, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const num_nodes = node_token_ids.size(); @@ -951,7 +951,7 @@ __device__ inline bool is_json_null_literal(char const* json, // Zero-row `List>>` for empty input. Sibling of `make_empty_map`; // the struct's value child is an empty `List` instead of an empty `STRING`. -std::unique_ptr make_empty_map_array(rmm::cuda_stream_view stream, +std::unique_ptr make_empty_map_array(cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return make_empty_map_from_value( @@ -1017,7 +1017,7 @@ struct element_classify_fn { std::unique_ptr from_json_to_raw_map(cudf::strings_column_view const& input, json_parse_options options, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -1090,7 +1090,7 @@ std::unique_ptr from_json_to_raw_map(cudf::strings_column_view con std::unique_ptr from_json_to_raw_map_array_values( cudf::strings_column_view const& input, json_parse_options options, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/from_json_to_raw_map_debug.cuh b/src/main/cpp/src/from_json_to_raw_map_debug.cuh index 481002c369..0eb7f09080 100644 --- a/src/main/cpp/src/from_json_to_raw_map_debug.cuh +++ b/src/main/cpp/src/from_json_to_raw_map_debug.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * Copyright (c) 2023-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -25,9 +25,10 @@ #include // -#include #include +#include + // #include @@ -61,7 +62,7 @@ template void print_debug(rmm::device_uvector const& input, std::string const& name, std::string const& separator, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const h_input = cudf::detail::make_host_vector_sync( cudf::device_span{input.data(), input.size()}, stream); @@ -78,7 +79,7 @@ void print_debug(rmm::device_uvector const& input, template void print_map_debug(rmm::device_uvector const& input, std::string const& name, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const h_input = cudf::detail::make_host_vector_sync( cudf::device_span{input.data(), input.size()}, stream); @@ -94,7 +95,7 @@ void print_map_debug(rmm::device_uvector const& input, template void print_pair_debug(rmm::device_uvector const& input, std::string const& name, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const h_input = cudf::detail::make_host_vector_sync( cudf::device_span{input.data(), input.size()}, stream); @@ -111,7 +112,7 @@ void print_pair_debug(rmm::device_uvector const& input, void print_output_spark_map(std::unique_ptr const& list_offsets, std::unique_ptr const& extracted_keys, std::unique_ptr const& extracted_values, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (extracted_keys->size() == 0) { std::cerr << "Extract keys-values are all empty.\n" << std::endl; diff --git a/src/main/cpp/src/from_json_to_structs.cu b/src/main/cpp/src/from_json_to_structs.cu index 30abeeeb2f..ceb0753fd2 100644 --- a/src/main/cpp/src/from_json_to_structs.cu +++ b/src/main/cpp/src/from_json_to_structs.cu @@ -35,7 +35,6 @@ #include #include -#include #include #include #include @@ -46,6 +45,7 @@ #include #include #include +#include #include #include #include @@ -160,7 +160,7 @@ std::pair generate_stru std::unique_ptr make_empty_column_from_schema( schema_element_with_precision const& schema, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (schema.type.id() == cudf::type_id::LIST) { @@ -191,7 +191,7 @@ std::unique_ptr make_empty_column_from_schema( void nullify_rows(cudf::column& input, std::span row_indices, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (row_indices.empty()) { return; } @@ -266,7 +266,7 @@ void nullify_rows(cudf::column& input, cudf::size_type null_count, rmm::device_buffer&& null_mask, bool did_nullify_schema_mismatch_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { std::vector> children; @@ -292,7 +292,7 @@ void nullify_rows(cudf::column& input, cudf::size_type null_count, rmm::device_buffer&& null_mask, bool did_nullify_schema_mismatch_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (did_nullify_schema_mismatch_rows && null_count > 0) { @@ -312,7 +312,7 @@ void nullify_rows(cudf::column& input, using string_index_pair = cuda::std::pair; std::unique_ptr cast_strings_to_booleans(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -367,7 +367,7 @@ std::unique_ptr cast_strings_to_booleans(cudf::column_view const& std::unique_ptr cast_strings_to_integers(cudf::column_view const& input, cudf::data_type output_type, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -435,7 +435,7 @@ std::unique_ptr cast_strings_to_integers(cudf::column_view const& } std::pair, bool> try_remove_quotes_for_floats( - cudf::column_view const& input, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + cudf::column_view const& input, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -521,7 +521,7 @@ std::pair, bool> try_remove_quotes_for_floats( std::unique_ptr cast_strings_to_floats(cudf::column_view const& input, cudf::data_type output_type, bool allow_nonnumeric_numbers, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -548,7 +548,7 @@ std::unique_ptr cast_strings_to_decimals(cudf::column_view const& cudf::data_type output_type, int precision, bool is_us_locale, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -597,7 +597,7 @@ std::unique_ptr cast_strings_to_decimals(cudf::column_view const& in_offsets + 1, plus_op, count_type{0, 0}, - stream.value()); + stream.get()); auto d_temp_storage = rmm::device_buffer{temp_storage_bytes, stream}; cub::DeviceSegmentedReduce::Reduce(d_temp_storage.data(), temp_storage_bytes, @@ -608,7 +608,7 @@ std::unique_ptr cast_strings_to_decimals(cudf::column_view const& in_offsets + 1, plus_op, count_type{0, 0}, - stream.value()); + stream.get()); } auto const out_size_it = spark_rapids_jni::util::make_counting_transform_iterator( @@ -691,7 +691,7 @@ std::unique_ptr cast_strings_to_decimals(cudf::column_view const& std::pair, bool> try_remove_quotes( cudf::strings_column_view const& input, bool nullify_if_not_quoted, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -776,7 +776,7 @@ std::unique_ptr convert_data_type(InputType&& input, bool allow_nonnumeric_numbers, bool is_us_locale, bool did_nullify_schema_mismatch_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -993,7 +993,7 @@ std::unique_ptr from_json_to_structs(cudf::strings_column_view con bool allow_nonnumeric_numbers, bool allow_unquoted_control, bool is_us_locale, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const [concat_input, delimiter, should_be_nullified] = @@ -1093,7 +1093,7 @@ std::unique_ptr from_json_to_structs(cudf::strings_column_view con bool allow_nonnumeric_numbers, bool allow_unquoted_control, bool is_us_locale, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -1120,7 +1120,7 @@ std::unique_ptr convert_from_strings(cudf::strings_column_view con std::vector const& precisions, bool allow_nonnumeric_numbers, bool is_us_locale, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -1146,7 +1146,7 @@ std::unique_ptr convert_from_strings(cudf::strings_column_view con std::unique_ptr remove_quotes(cudf::strings_column_view const& input, bool nullify_if_not_quoted, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/get_json_object.cu b/src/main/cpp/src/get_json_object.cu index ae6b380e46..4907b9f4a6 100644 --- a/src/main/cpp/src/get_json_object.cu +++ b/src/main/cpp/src/get_json_object.cu @@ -870,7 +870,7 @@ struct kernel_launcher { static void exec(cudf::column_device_view const& input, cudf::device_span path_data, int8_t* max_path_depth_exceeded, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // The optimal values for block_size and min_block_per_sm were found through testing, // which are either 128-8 or 256-4. The pair 128-8 seems a bit better. @@ -885,7 +885,7 @@ struct kernel_launcher { auto const num_blocks = cudf::util::div_rounding_up_safe(num_threads_per_row * input.size(), static_cast(block_size)); get_json_object_kernel - <<>>( + <<>>( input, path_data, num_threads_per_row, max_path_depth_exceeded); } }; @@ -906,7 +906,7 @@ std::tuple>, construct_path_commands( std::vector const>> const& json_paths, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // Concatenate all names from path instructions. auto h_inst_names = [&] { @@ -966,7 +966,7 @@ construct_path_commands( int64_t calc_scratch_size(cudf::strings_column_view const& input, cudf::detail::input_offsetalator const& in_offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const max_row_size = thrust::transform_reduce( rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), @@ -1018,7 +1018,7 @@ std::vector> get_json_object_batch( std::vector const>> const& json_paths, int64_t scratch_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const [d_json_paths, h_json_paths, d_inst_names, h_inst_names] = @@ -1173,7 +1173,7 @@ std::vector> get_json_object( json_paths, int64_t memory_budget_bytes, int32_t parallel_override, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const num_outputs = json_paths.size(); @@ -1247,7 +1247,7 @@ std::vector> get_json_object( std::unique_ptr get_json_object( cudf::strings_column_view const& input, std::vector> const& instructions, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -1260,7 +1260,7 @@ std::vector> get_json_object_multiple_paths( json_paths, int64_t memory_budget_bytes, int32_t parallel_override, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/get_json_object.hpp b/src/main/cpp/src/get_json_object.hpp index dbf62de5c2..d4b856971d 100644 --- a/src/main/cpp/src/get_json_object.hpp +++ b/src/main/cpp/src/get_json_object.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -20,6 +20,8 @@ #include +#include + #include #include @@ -44,7 +46,7 @@ enum class path_instruction_type : int8_t { WILDCARD, INDEX, NAMED }; std::unique_ptr get_json_object( cudf::strings_column_view const& input, std::vector> const& instructions, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -66,7 +68,7 @@ std::vector> get_json_object_multiple_paths( json_paths, int64_t memory_budget_bytes, int32_t parallel_override, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/hash/hash.hpp b/src/main/cpp/src/hash/hash.hpp index 1dd46c4cc7..cd978ffde7 100644 --- a/src/main/cpp/src/hash/hash.hpp +++ b/src/main/cpp/src/hash/hash.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -19,9 +19,10 @@ #include #include -#include #include +#include + namespace spark_rapids_jni { constexpr int64_t DEFAULT_XXHASH64_SEED = 42; @@ -40,7 +41,7 @@ constexpr int MAX_STACK_DEPTH = 8; std::unique_ptr murmur_hash3_32( cudf::table_view const& input, uint32_t seed = 0, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -56,7 +57,7 @@ std::unique_ptr murmur_hash3_32( std::unique_ptr xxhash64( cudf::table_view const& input, int64_t seed = DEFAULT_XXHASH64_SEED, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -70,7 +71,7 @@ std::unique_ptr xxhash64( */ std::unique_ptr hive_hash( cudf::table_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -85,7 +86,7 @@ std::unique_ptr hive_hash( */ std::unique_ptr sha224_nulls_preserved( cudf::column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -100,7 +101,7 @@ std::unique_ptr sha224_nulls_preserved( */ std::unique_ptr sha256_nulls_preserved( cudf::column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -115,7 +116,7 @@ std::unique_ptr sha256_nulls_preserved( */ std::unique_ptr sha384_nulls_preserved( cudf::column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -130,7 +131,7 @@ std::unique_ptr sha384_nulls_preserved( */ std::unique_ptr sha512_nulls_preserved( cudf::column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/hash/hive_hash.cu b/src/main/cpp/src/hash/hive_hash.cu index 9b51e770d0..757fbc42c9 100644 --- a/src/main/cpp/src/hash/hive_hash.cu +++ b/src/main/cpp/src/hash/hive_hash.cu @@ -23,12 +23,12 @@ #include #include -#include #include #include #include #include +#include #include #include @@ -471,7 +471,7 @@ void check_nested_depth(cudf::table_view const& input) } // namespace std::unique_ptr hive_hash(cudf::table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto output = cudf::make_numeric_column(cudf::data_type(cudf::type_to_id()), diff --git a/src/main/cpp/src/hash/murmur_hash.cu b/src/main/cpp/src/hash/murmur_hash.cu index 66043eb536..3df6370142 100644 --- a/src/main/cpp/src/hash/murmur_hash.cu +++ b/src/main/cpp/src/hash/murmur_hash.cu @@ -22,11 +22,11 @@ #include #include -#include #include #include #include +#include #include #include @@ -191,7 +191,7 @@ void check_hash_compatibility(cudf::table_view const& input) std::unique_ptr murmur_hash3_32(cudf::table_view const& input, uint32_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto output = diff --git a/src/main/cpp/src/hash/sha.cpp b/src/main/cpp/src/hash/sha.cpp index 6b8551a5db..f82bb88869 100644 --- a/src/main/cpp/src/hash/sha.cpp +++ b/src/main/cpp/src/hash/sha.cpp @@ -19,19 +19,20 @@ #include #include -#include #include +#include + #include namespace { using HashFunction = std::unique_ptr (*)(cudf::table_view const&, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref); std::unique_ptr sha_impl(HashFunction hash_function, cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.is_empty()) { return cudf::make_empty_column(cudf::data_type{cudf::type_id::STRING}); } @@ -55,28 +56,28 @@ std::unique_ptr sha_impl(HashFunction hash_function, namespace spark_rapids_jni { std::unique_ptr sha224_nulls_preserved(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return sha_impl(cudf::hashing::sha224, input, stream, mr); } std::unique_ptr sha256_nulls_preserved(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return sha_impl(cudf::hashing::sha256, input, stream, mr); } std::unique_ptr sha384_nulls_preserved(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return sha_impl(cudf::hashing::sha384, input, stream, mr); } std::unique_ptr sha512_nulls_preserved(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return sha_impl(cudf::hashing::sha512, input, stream, mr); diff --git a/src/main/cpp/src/hash/xxhash64.cu b/src/main/cpp/src/hash/xxhash64.cu index 89645ac971..eebc3fc4fb 100644 --- a/src/main/cpp/src/hash/xxhash64.cu +++ b/src/main/cpp/src/hash/xxhash64.cu @@ -22,13 +22,13 @@ #include #include -#include #include #include #include #include #include +#include #include namespace spark_rapids_jni { @@ -552,7 +552,7 @@ void check_nested_depth(cudf::table_view const& input) std::unique_ptr xxhash64(cudf::table_view const& input, int64_t _seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { hash_value_type seed = static_cast(_seed); diff --git a/src/main/cpp/src/hex.cu b/src/main/cpp/src/hex.cu index a24e531f8f..cd48171443 100644 --- a/src/main/cpp/src/hex.cu +++ b/src/main/cpp/src/hex.cu @@ -25,10 +25,10 @@ #include #include -#include #include #include +#include #include #include #include @@ -83,7 +83,7 @@ struct write_hex_fn { } // namespace std::unique_ptr bytes_to_hex(cudf::strings_column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.is_empty()) { return cudf::make_empty_column(cudf::type_id::STRING); } @@ -119,7 +119,7 @@ std::unique_ptr bytes_to_hex(cudf::strings_column_view const& inpu } // namespace detail std::unique_ptr bytes_to_hex(cudf::strings_column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/histogram.cu b/src/main/cpp/src/histogram.cu index 0467430dcd..fbc3d064ed 100644 --- a/src/main/cpp/src/histogram.cu +++ b/src/main/cpp/src/histogram.cu @@ -164,7 +164,7 @@ struct percentile_dispatcher { cudf::device_span percentages, bool has_null, cudf::size_type num_histograms, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { // Returns all nulls for totally empty input. @@ -255,7 +255,7 @@ std::unique_ptr wrap_in_list(std::unique_ptr&& input cudf::size_type null_count, cudf::size_type num_histograms, cudf::size_type num_percentages, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input->size() == 0) { return cudf::make_empty_lists_column(input->type()); } @@ -275,7 +275,7 @@ std::unique_ptr wrap_in_list(std::unique_ptr&& input std::unique_ptr create_histogram_if_valid(cudf::column_view const& values, cudf::column_view const& frequencies, bool output_as_lists, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS( @@ -416,7 +416,7 @@ std::unique_ptr create_histogram_if_valid(cudf::column_view const& std::unique_ptr percentile_from_histogram(cudf::column_view const& input, std::vector const& percentages, bool output_as_list, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { check_input(input, percentages); diff --git a/src/main/cpp/src/histogram.hpp b/src/main/cpp/src/histogram.hpp index e9ac759197..03e51db407 100644 --- a/src/main/cpp/src/histogram.hpp +++ b/src/main/cpp/src/histogram.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * Copyright (c) 2023-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -19,9 +19,10 @@ #include // -#include #include +#include + namespace spark_rapids_jni { /** @@ -51,7 +52,7 @@ std::unique_ptr create_histogram_if_valid( cudf::column_view const& values, cudf::column_view const& frequencies, bool output_as_lists, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -71,7 +72,7 @@ std::unique_ptr percentile_from_histogram( cudf::column_view const& input, std::vector const& percentage, bool output_as_lists, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/hyper_log_log_plus_plus.hpp b/src/main/cpp/src/hyper_log_log_plus_plus.hpp index 5b6b799ca9..93ba116bef 100644 --- a/src/main/cpp/src/hyper_log_log_plus_plus.hpp +++ b/src/main/cpp/src/hyper_log_log_plus_plus.hpp @@ -20,7 +20,7 @@ #include #include -#include +#include namespace spark_rapids_jni { diff --git a/src/main/cpp/src/iceberg/iceberg_bucket.cu b/src/main/cpp/src/iceberg/iceberg_bucket.cu index d2bcfe0849..f486f8e785 100644 --- a/src/main/cpp/src/iceberg/iceberg_bucket.cu +++ b/src/main/cpp/src/iceberg/iceberg_bucket.cu @@ -383,7 +383,7 @@ struct bucket_decimal128_fn { template void generate_buckets(GeneratorFunc generator, cudf::mutable_column_view output, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { thrust::tabulate(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), output.begin(), @@ -393,7 +393,7 @@ void generate_buckets(GeneratorFunc generator, std::unique_ptr compute_bucket_impl(cudf::column_view const& input, int32_t num_buckets, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(num_buckets > 0, "num_buckets must be positive"); @@ -467,7 +467,7 @@ std::unique_ptr compute_bucket_impl(cudf::column_view const& input std::unique_ptr compute_bucket(cudf::column_view const& input, int32_t num_buckets, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/iceberg/iceberg_bucket.hpp b/src/main/cpp/src/iceberg/iceberg_bucket.hpp index 014f23ed86..84616746ff 100644 --- a/src/main/cpp/src/iceberg/iceberg_bucket.hpp +++ b/src/main/cpp/src/iceberg/iceberg_bucket.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2025, NVIDIA CORPORATION. + * Copyright (c) 2025-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -21,6 +21,8 @@ #include #include +#include + namespace spark_rapids_jni { /** @@ -50,7 +52,7 @@ namespace spark_rapids_jni { std::unique_ptr compute_bucket( cudf::column_view const& input, int32_t num_buckets, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/iceberg/iceberg_datetime_util.cu b/src/main/cpp/src/iceberg/iceberg_datetime_util.cu index 00bc35f057..ff0a213216 100644 --- a/src/main/cpp/src/iceberg/iceberg_datetime_util.cu +++ b/src/main/cpp/src/iceberg/iceberg_datetime_util.cu @@ -136,7 +136,7 @@ struct hours_from_epoch_for_ts_fn { }; std::unique_ptr compute_years_from_epoch(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.is_empty()) { return cudf::make_empty_column(cudf::data_type{cudf::type_id::INT32}); } @@ -165,7 +165,7 @@ std::unique_ptr compute_years_from_epoch(cudf::column_view const& } std::unique_ptr compute_months_from_epoch(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.is_empty()) { return cudf::make_empty_column(cudf::data_type{cudf::type_id::INT32}); } @@ -194,7 +194,7 @@ std::unique_ptr compute_months_from_epoch(cudf::column_view const& } std::unique_ptr compute_days_from_epoch(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.is_empty()) { @@ -225,7 +225,7 @@ std::unique_ptr compute_days_from_epoch(cudf::column_view const& i } std::unique_ptr compute_hours_from_epoch(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.is_empty()) { return cudf::make_empty_column(cudf::data_type{cudf::type_id::INT32}); } @@ -251,7 +251,7 @@ std::unique_ptr compute_hours_from_epoch(cudf::column_view const& } // anonymous namespace std::unique_ptr years_from_epoch(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -259,7 +259,7 @@ std::unique_ptr years_from_epoch(cudf::column_view const& input, } std::unique_ptr months_from_epoch(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -267,7 +267,7 @@ std::unique_ptr months_from_epoch(cudf::column_view const& input, } std::unique_ptr days_from_epoch(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -275,7 +275,7 @@ std::unique_ptr days_from_epoch(cudf::column_view const& input, } std::unique_ptr hours_from_epoch(cudf::column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/iceberg/iceberg_datetime_util.hpp b/src/main/cpp/src/iceberg/iceberg_datetime_util.hpp index cb401012de..29bb103fd4 100644 --- a/src/main/cpp/src/iceberg/iceberg_datetime_util.hpp +++ b/src/main/cpp/src/iceberg/iceberg_datetime_util.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2025, NVIDIA CORPORATION. + * Copyright (c) 2025-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -21,6 +21,8 @@ #include #include +#include + #include namespace spark_rapids_jni { @@ -37,7 +39,7 @@ namespace spark_rapids_jni { */ std::unique_ptr years_from_epoch( cudf::column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -52,7 +54,7 @@ std::unique_ptr years_from_epoch( */ std::unique_ptr months_from_epoch( cudf::column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -67,7 +69,7 @@ std::unique_ptr months_from_epoch( */ std::unique_ptr days_from_epoch( cudf::column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -82,7 +84,7 @@ std::unique_ptr days_from_epoch( */ std::unique_ptr hours_from_epoch( cudf::column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/iceberg/iceberg_truncate.cu b/src/main/cpp/src/iceberg/iceberg_truncate.cu index 9809fb4b47..eff3e9bc10 100644 --- a/src/main/cpp/src/iceberg/iceberg_truncate.cu +++ b/src/main/cpp/src/iceberg/iceberg_truncate.cu @@ -131,7 +131,7 @@ template void truncate_integral_and_fill(std::unique_ptr& output, cudf::column_device_view d_input, int32_t width, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { thrust::tabulate(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), output->mutable_view().begin(), @@ -141,7 +141,7 @@ void truncate_integral_and_fill(std::unique_ptr& output, std::unique_ptr truncate_integral_impl(cudf::column_view const& input, int32_t width, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(width != 0, "Width must not be zero"); @@ -170,7 +170,7 @@ std::unique_ptr truncate_integral_impl(cudf::column_view const& in std::unique_ptr truncate_string_impl(cudf::column_view const& input, int32_t truncate_length, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input.type().id() == cudf::type_id::STRING, "Input must be STRING"); @@ -198,7 +198,7 @@ std::unique_ptr truncate_string_impl(cudf::column_view const& inpu std::unique_ptr truncate_binary_impl(cudf::column_view const& input, int32_t truncate_length, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(truncate_length > 0, "Length must be positive"); @@ -243,7 +243,7 @@ std::unique_ptr truncate_binary_impl(cudf::column_view const& inpu std::unique_ptr truncate_integral(cudf::column_view const& input, int32_t width, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -252,7 +252,7 @@ std::unique_ptr truncate_integral(cudf::column_view const& input, std::unique_ptr truncate_string(cudf::column_view const& input, int32_t truncate_length, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -261,7 +261,7 @@ std::unique_ptr truncate_string(cudf::column_view const& input, std::unique_ptr truncate_binary(cudf::column_view const& input, int32_t truncate_length, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/iceberg/iceberg_truncate.hpp b/src/main/cpp/src/iceberg/iceberg_truncate.hpp index 3a80936cc2..4896039f2a 100644 --- a/src/main/cpp/src/iceberg/iceberg_truncate.hpp +++ b/src/main/cpp/src/iceberg/iceberg_truncate.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2025, NVIDIA CORPORATION. + * Copyright (c) 2025-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -21,6 +21,8 @@ #include #include +#include + namespace spark_rapids_jni { /** @@ -54,7 +56,7 @@ namespace spark_rapids_jni { std::unique_ptr truncate_integral( cudf::column_view const& input, int32_t width, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -77,7 +79,7 @@ std::unique_ptr truncate_integral( std::unique_ptr truncate_string( cudf::column_view const& input, int32_t length, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -93,7 +95,7 @@ std::unique_ptr truncate_string( std::unique_ptr truncate_binary( cudf::column_view const& input, int32_t length, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/join_primitives.cu b/src/main/cpp/src/join_primitives.cu index e37b237891..abcde659c5 100644 --- a/src/main/cpp/src/join_primitives.cu +++ b/src/main/cpp/src/join_primitives.cu @@ -31,10 +31,10 @@ #include #include -#include #include #include +#include #include #include #include @@ -96,7 +96,7 @@ filter_by_conditional_impl(cudf::device_span left_indices cudf::table_device_view const& right_table, cudf::ast::detail::expression_device_view device_expression_data, cudf::size_type num_intermediates, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const num_pairs = left_indices.size(); @@ -135,13 +135,13 @@ filter_by_conditional_impl(cudf::device_span left_indices auto const shmem_size = static_cast(block_size) * static_cast(per_thread_bytes); filter_join_indices_kernel - <<>>(left_indices.data(), - right_indices.data(), - num_pairs, - left_table, - right_table, - device_expression_data, - keep_mask.data()); + <<>>(left_indices.data(), + right_indices.data(), + num_pairs, + left_table, + right_table, + device_expression_data, + keep_mask.data()); // Surface any kernel launch errors immediately CUDF_CUDA_TRY(cudaPeekAtLastError()); @@ -183,7 +183,7 @@ sort_merge_inner_join(cudf::table_view const& left_keys, cudf::sorted is_left_sorted, cudf::sorted is_right_sorted, cudf::null_equality compare_nulls, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -208,7 +208,7 @@ std::pair, rmm::device_uvector left_indices, cudf::table_view const& left_table, cudf::table_view const& right_table, cudf::ast::expression const& binary_predicate, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -308,13 +308,13 @@ namespace { std::pair, cudf::size_type> compute_side_match_info( cudf::device_span indices, cudf::size_type table_size, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // Create a boolean mask to track which rows have matches // Note: Temporary buffers use current device resource for allocation auto has_match = rmm::device_uvector(table_size, stream, cudf::get_current_device_resource_ref()); - CUDF_CUDA_TRY(cudaMemsetAsync(has_match.data(), 0, has_match.size(), stream.value())); + CUDF_CUDA_TRY(cudaMemsetAsync(has_match.data(), 0, has_match.size(), stream.get())); // Mark rows that have matches thrust::for_each(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), @@ -343,7 +343,7 @@ void populate_outer_result(cudf::size_type table_size, cudf::size_type unmatched_offset, rmm::device_uvector& out_indices, rmm::device_uvector& out_other_indices, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // Copy unmatched rows auto unmatched_iter = out_indices.begin() + unmatched_offset; @@ -369,7 +369,7 @@ make_left_outer(cudf::device_span left_indices, cudf::device_span right_indices, cudf::size_type left_table_size, cudf::size_type right_table_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -414,7 +414,7 @@ make_full_outer(cudf::device_span left_indices, cudf::device_span right_indices, cudf::size_type left_table_size, cudf::size_type right_table_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -476,7 +476,7 @@ make_full_outer(cudf::device_span left_indices, rmm::device_uvector make_semi( cudf::device_span left_indices, cudf::size_type left_table_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -520,7 +520,7 @@ rmm::device_uvector make_semi( rmm::device_uvector make_anti( cudf::device_span left_indices, cudf::size_type left_table_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -567,7 +567,7 @@ rmm::device_uvector make_anti( std::unique_ptr get_matched_rows(cudf::device_span gather_map, cudf::size_type table_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -580,7 +580,7 @@ std::unique_ptr get_matched_rows(cudf::device_span(); // Initialize all to false - CUDF_CUDA_TRY(cudaMemsetAsync(result_data, 0, table_size, stream.value())); + CUDF_CUDA_TRY(cudaMemsetAsync(result_data, 0, table_size, stream.get())); // Mark rows that appear in the gather map as true thrust::for_each(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), diff --git a/src/main/cpp/src/join_primitives.hpp b/src/main/cpp/src/join_primitives.hpp index f3476b01ab..a10589ba05 100644 --- a/src/main/cpp/src/join_primitives.hpp +++ b/src/main/cpp/src/join_primitives.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2025, NVIDIA CORPORATION. + * Copyright (c) 2025-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -24,9 +24,10 @@ #include #include -#include #include +#include + #include #include @@ -67,7 +68,7 @@ sort_merge_inner_join(cudf::table_view const& left_keys, cudf::sorted is_left_sorted = cudf::sorted::NO, cudf::sorted is_right_sorted = cudf::sorted::NO, cudf::null_equality compare_nulls = cudf::null_equality::EQUAL, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -88,7 +89,7 @@ std::pair, rmm::device_uvector left_indices, cudf::device_span right_indices, cudf::size_type left_table_size, cudf::size_type right_table_size, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -170,7 +171,7 @@ make_full_outer(cudf::device_span left_indices, cudf::device_span right_indices, cudf::size_type left_table_size, cudf::size_type right_table_size, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); // ============================================================================= @@ -194,7 +195,7 @@ make_full_outer(cudf::device_span left_indices, rmm::device_uvector make_semi( cudf::device_span indices, cudf::size_type table_size, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -213,7 +214,7 @@ rmm::device_uvector make_semi( rmm::device_uvector make_anti( cudf::device_span indices, cudf::size_type table_size, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); // ============================================================================= @@ -237,7 +238,7 @@ rmm::device_uvector make_anti( std::unique_ptr get_matched_rows( cudf::device_span gather_map, cudf::size_type table_size, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/json_utils.cu b/src/main/cpp/src/json_utils.cu index baf0982410..e2698ad66d 100644 --- a/src/main/cpp/src/json_utils.cu +++ b/src/main/cpp/src/json_utils.cu @@ -24,7 +24,6 @@ #include #include -#include #include #include @@ -32,6 +31,7 @@ #include #include #include +#include #include #include #include @@ -119,7 +119,7 @@ __host__ __device__ constexpr std::uint8_t delimiter_candidate(int candidate_ind std::tuple, char, std::unique_ptr> concat_json( cudf::strings_column_view const& input, bool nullify_invalid_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.is_empty()) { @@ -209,7 +209,7 @@ std::tuple, char, std::unique_ptr, char, std::unique_ptr, char, std::unique_ptr, char, std::unique_ptr> concat_json( cudf::strings_column_view const& input, bool nullify_invalid_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/json_utils.hpp b/src/main/cpp/src/json_utils.hpp index 33ace7f259..9a5a2951d5 100644 --- a/src/main/cpp/src/json_utils.hpp +++ b/src/main/cpp/src/json_utils.hpp @@ -20,9 +20,10 @@ #include #include -#include #include +#include + #include namespace spark_rapids_jni { @@ -43,7 +44,7 @@ struct json_parse_options { std::unique_ptr from_json_to_raw_map( cudf::strings_column_view const& input, json_parse_options options, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -65,7 +66,7 @@ std::unique_ptr from_json_to_raw_map( std::unique_ptr from_json_to_raw_map_array_values( cudf::strings_column_view const& input, json_parse_options options, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -85,7 +86,7 @@ std::unique_ptr from_json_to_structs( bool allow_nonnumeric_numbers, bool allow_unquoted_control, bool is_us_locale, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -101,7 +102,7 @@ std::unique_ptr convert_from_strings( std::vector const& precisions, bool allow_nonnumeric_numbers, bool is_us_locale, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -113,7 +114,7 @@ std::unique_ptr convert_from_strings( std::unique_ptr remove_quotes( cudf::strings_column_view const& input, bool nullify_if_not_quoted, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -139,7 +140,7 @@ std::unique_ptr remove_quotes( std::tuple, char, std::unique_ptr> concat_json( cudf::strings_column_view const& input, bool nullify_invalid_rows = false, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/list_slice.cu b/src/main/cpp/src/list_slice.cu index 2d08b8600b..94b8d54582 100644 --- a/src/main/cpp/src/list_slice.cu +++ b/src/main/cpp/src/list_slice.cu @@ -26,10 +26,10 @@ #include #include -#include #include #include +#include #include using namespace cudf; @@ -40,7 +40,7 @@ namespace detail { namespace { -void assert_start_is_not_zero(column_device_view const& start, rmm::cuda_stream_view stream) +void assert_start_is_not_zero(column_device_view const& start, cuda::stream_ref stream) { bool start_valid = thrust::all_of(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), @@ -53,7 +53,7 @@ void assert_start_is_not_zero(column_device_view const& start, rmm::cuda_stream_ CUDF_EXPECTS(start_valid, "Invalid start value: start must not be 0"); } -void assert_length_is_not_negative(column_device_view const& length, rmm::cuda_stream_view stream) +void assert_length_is_not_negative(column_device_view const& length, cuda::stream_ref stream) { bool length_valid = thrust::all_of(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), @@ -148,7 +148,7 @@ auto generate_starts_and_sizes(size_type const* offsets_of_input_lists, size_type const num_rows, SRART_ITERATOR const start_iterator, LENGTH_ITERATOR const length_iterator, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto starts = make_numeric_column(data_type{type_id::INT32}, num_rows, @@ -162,7 +162,7 @@ auto generate_starts_and_sizes(size_type const* offsets_of_input_lists, cudf::get_current_device_resource_ref()); constexpr int block_size = 256; auto grid = cudf::detail::grid_1d{num_rows, block_size}; - compute_starts_and_sizes_kernel<<>>( + compute_starts_and_sizes_kernel<<>>( offsets_of_input_lists, num_rows, start_iterator, @@ -175,7 +175,7 @@ auto generate_starts_and_sizes(size_type const* offsets_of_input_lists, std::unique_ptr legal_list_slice(lists_column_view const& input, column_view const& starts, column_view const& sizes, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const num_rows = input.size(); @@ -187,7 +187,7 @@ std::unique_ptr legal_list_slice(lists_column_view const& input, constexpr int block_size = 256; auto grid = cudf::detail::grid_1d{num_rows, block_size}; rmm::device_uvector gather_map(num_total_elements, stream); - compute_gather_map<<>>( + compute_gather_map<<>>( num_rows, input.offsets_begin(), starts.begin(), @@ -220,7 +220,7 @@ std::unique_ptr list_slice(lists_column_view const& input, size_type const start, size_type const length, bool check_start_length, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (check_start_length) { @@ -244,7 +244,7 @@ std::unique_ptr list_slice(lists_column_view const& input, size_type const start, column_view const& length, bool check_start_length, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(length.type().id() == type_id::INT32, @@ -277,7 +277,7 @@ std::unique_ptr list_slice(lists_column_view const& input, column_view const& start, size_type const length, bool check_start_length, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(start.type().id() == type_id::INT32, @@ -310,7 +310,7 @@ std::unique_ptr list_slice(lists_column_view const& input, column_view const& start, column_view const& length, bool check_start_length, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(start.type().id() == type_id::INT32, @@ -353,7 +353,7 @@ std::unique_ptr list_slice(lists_column_view const& input, size_type const start, size_type const length, bool check_start_length, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -364,7 +364,7 @@ std::unique_ptr list_slice(lists_column_view const& input, size_type const start, column_view const& length, bool check_start_length, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -375,7 +375,7 @@ std::unique_ptr list_slice(lists_column_view const& input, column_view const& start, size_type const length, bool check_start_length, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -386,7 +386,7 @@ std::unique_ptr list_slice(lists_column_view const& input, column_view const& start, column_view const& length, bool check_start_length, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/list_slice.hpp b/src/main/cpp/src/list_slice.hpp index 8daad1d266..e60ff02bd9 100644 --- a/src/main/cpp/src/list_slice.hpp +++ b/src/main/cpp/src/list_slice.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2025, NVIDIA CORPORATION. + * Copyright (c) 2025-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -21,6 +21,8 @@ #include #include +#include + namespace spark_rapids_jni { /** @@ -70,7 +72,7 @@ std::unique_ptr list_slice( cudf::size_type const start, cudf::size_type const length, bool check_start_length = true, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -123,7 +125,7 @@ std::unique_ptr list_slice( cudf::size_type const start, cudf::column_view const& length, bool check_start_length = true, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -176,7 +178,7 @@ std::unique_ptr list_slice( cudf::column_view const& start, cudf::size_type const length, bool check_start_length = true, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -239,7 +241,7 @@ std::unique_ptr list_slice( cudf::column_view const& start, cudf::column_view const& length, bool check_start_length = true, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/map.cu b/src/main/cpp/src/map.cu index ae5aa282b9..2efe5c595a 100644 --- a/src/main/cpp/src/map.cu +++ b/src/main/cpp/src/map.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2025, NVIDIA CORPORATION. + * Copyright (c) 2025-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -28,7 +28,7 @@ namespace spark_rapids_jni { std::unique_ptr sort_map_column(cudf::column_view const& input, cudf::order sort_order, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input.type().id() == cudf::type_id::LIST, diff --git a/src/main/cpp/src/map.hpp b/src/main/cpp/src/map.hpp index d8af75ecdd..4da344db0e 100644 --- a/src/main/cpp/src/map.hpp +++ b/src/main/cpp/src/map.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2025, NVIDIA CORPORATION. + * Copyright (c) 2025-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -18,7 +18,7 @@ #include #include -#include +#include namespace spark_rapids_jni { @@ -39,7 +39,7 @@ namespace spark_rapids_jni { std::unique_ptr sort_map_column( cudf::column_view const& input, cudf::order sort_order, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/map_utils.cu b/src/main/cpp/src/map_utils.cu index 2de8aed900..0eed231db3 100644 --- a/src/main/cpp/src/map_utils.cu +++ b/src/main/cpp/src/map_utils.cu @@ -32,7 +32,6 @@ #include #include -#include #include #include #include @@ -40,6 +39,7 @@ #include #include +#include #include #include @@ -174,7 +174,7 @@ struct phase1_state_summary { phase1_state_summary run_phase1_state(cudf::column_view const& input, bool throw_on_null_key, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) { auto const num_rows = input.size(); @@ -189,15 +189,15 @@ phase1_state_summary run_phase1_state(cudf::column_view const& input, { constexpr int block_size = 256; auto const grid_size = cudf::util::div_rounding_up_safe(num_rows, block_size); - compute_row_state_kernel<<>>(input.null_mask(), - lists_cv.offsets_begin(), - structs.null_mask(), - keys.null_mask(), - throw_on_null_key, - num_rows, - row_state.data(), - row_size.data()); - CUDF_CHECK_CUDA(stream.value()); + compute_row_state_kernel<<>>(input.null_mask(), + lists_cv.offsets_begin(), + structs.null_mask(), + keys.null_mask(), + throw_on_null_key, + num_rows, + row_state.data(), + row_size.data()); + CUDF_CHECK_CUDA(stream.get()); } // Max + Min reductions on row_state — the state ordering encodes both "must throw?" and @@ -207,28 +207,28 @@ phase1_state_summary run_phase1_state(cudf::column_view const& input, { std::size_t bytes = 0; CUDF_CUDA_TRY(cub::DeviceReduce::Max( - nullptr, bytes, row_state.data(), max_state_d.data(), num_rows, stream.value())); + nullptr, bytes, row_state.data(), max_state_d.data(), num_rows, stream.get())); rmm::device_buffer tmp(bytes, stream, temp_mr); CUDF_CUDA_TRY(cub::DeviceReduce::Max( - tmp.data(), bytes, row_state.data(), max_state_d.data(), num_rows, stream.value())); + tmp.data(), bytes, row_state.data(), max_state_d.data(), num_rows, stream.get())); } { std::size_t bytes = 0; CUDF_CUDA_TRY(cub::DeviceReduce::Min( - nullptr, bytes, row_state.data(), min_state_d.data(), num_rows, stream.value())); + nullptr, bytes, row_state.data(), min_state_d.data(), num_rows, stream.get())); rmm::device_buffer tmp(bytes, stream, temp_mr); CUDF_CUDA_TRY(cub::DeviceReduce::Min( - tmp.data(), bytes, row_state.data(), min_state_d.data(), num_rows, stream.value())); + tmp.data(), bytes, row_state.data(), min_state_d.data(), num_rows, stream.get())); } // Bundled D→H pull — two async copies, one stream sync. std::uint8_t max_state{}; std::uint8_t min_state{}; CUDF_CUDA_TRY(cudaMemcpyAsync( - &max_state, max_state_d.data(), sizeof(std::uint8_t), cudaMemcpyDefault, stream.value())); + &max_state, max_state_d.data(), sizeof(std::uint8_t), cudaMemcpyDefault, stream.get())); CUDF_CUDA_TRY(cudaMemcpyAsync( - &min_state, min_state_d.data(), sizeof(std::uint8_t), cudaMemcpyDefault, stream.value())); - stream.synchronize(); + &min_state, min_state_d.data(), sizeof(std::uint8_t), cudaMemcpyDefault, stream.get())); + stream.sync(); return phase1_state_summary{std::move(row_state), std::move(row_size), max_state, min_state}; } @@ -237,7 +237,7 @@ phase1_state_summary run_phase1_state(cudf::column_view const& input, bool is_valid_map(cudf::column_view const& input, bool throw_on_null_key, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -257,7 +257,7 @@ bool is_valid_map(cudf::column_view const& input, std::unique_ptr map_from_entries(cudf::column_view const& input, bool throw_on_null_key, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -284,7 +284,7 @@ std::unique_ptr map_from_entries(cudf::column_view const& input, auto const structs = lists_cv.child(); rmm::device_uvector out_offsets(num_rows + 1, stream, mr); - CUDF_CUDA_TRY(cudaMemsetAsync(out_offsets.data(), 0, sizeof(cudf::size_type), stream.value())); + CUDF_CUDA_TRY(cudaMemsetAsync(out_offsets.data(), 0, sizeof(cudf::size_type), stream.get())); thrust::inclusive_scan(rmm::exec_policy_nosync(stream, temp_mr), p1.row_size.begin(), p1.row_size.end(), @@ -295,8 +295,8 @@ std::unique_ptr map_from_entries(cudf::column_view const& input, out_offsets.data() + num_rows, sizeof(cudf::size_type), cudaMemcpyDefault, - stream.value())); - stream.synchronize(); + stream.get())); + stream.sync(); // ── Phase 2: clean output construction (no dirty intermediate result) ───── // 2a. Null mask from row_state via an explicit `state == STATE_VALID` predicate. Avoids @@ -317,12 +317,12 @@ std::unique_ptr map_from_entries(cudf::column_view const& input, { constexpr int block_size = 256; auto const grid_size = cudf::util::div_rounding_up_safe(num_rows, block_size); - build_gather_map_kernel<<>>(p1.row_state.data(), - lists_cv.offsets_begin(), - out_offsets.data(), - num_rows, - gather_map.data()); - CUDF_CHECK_CUDA(stream.value()); + build_gather_map_kernel<<>>(p1.row_state.data(), + lists_cv.offsets_begin(), + out_offsets.data(), + num_rows, + gather_map.data()); + CUDF_CHECK_CUDA(stream.get()); } // 2c. Single gather over the struct child — handles arbitrary nested key/value types. diff --git a/src/main/cpp/src/map_utils.hpp b/src/main/cpp/src/map_utils.hpp index 497d7a76b5..50b28f0f9d 100644 --- a/src/main/cpp/src/map_utils.hpp +++ b/src/main/cpp/src/map_utils.hpp @@ -20,7 +20,7 @@ #include #include -#include +#include namespace spark_rapids_jni { @@ -58,7 +58,7 @@ namespace spark_rapids_jni { [[nodiscard]] bool is_valid_map( cudf::column_view const& input, bool throw_on_null_key, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -97,7 +97,7 @@ namespace spark_rapids_jni { [[nodiscard]] std::unique_ptr map_from_entries( cudf::column_view const& input, bool throw_on_null_key, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/map_zip_with_utils.cu b/src/main/cpp/src/map_zip_with_utils.cu index 6f158fee10..db3e07ca6f 100644 --- a/src/main/cpp/src/map_zip_with_utils.cu +++ b/src/main/cpp/src/map_zip_with_utils.cu @@ -50,7 +50,7 @@ namespace { std::unique_ptr generate_labels( lists_column_view const& input, size_type n_elements, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()) { auto labels = make_numeric_column( @@ -112,8 +112,7 @@ std::unique_ptr generate_labels( std::unique_ptr indices_of( lists_column_view const& search_keys, // Column containing lists of keys to search for lists_column_view const& search_values, // Column containing lists of values to search through - rmm::cuda_stream_view stream = - cudf::get_default_stream(), // CUDA stream for asynchronous execution + cuda::stream_ref stream = cudf::get_default_stream(), // CUDA stream for asynchronous execution rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()) // Memory resource for allocations { @@ -310,7 +309,7 @@ std::unique_ptr indices_of( std::unique_ptr map_zip( cudf::lists_column_view const& col1, // First map column containing key-value pairs cudf::lists_column_view const& col2, // Second map column containing key-value pairs - rmm::cuda_stream_view stream, // CUDA stream for asynchronous execution + cuda::stream_ref stream, // CUDA stream for asynchronous execution rmm::device_async_resource_ref mr) // Memory resource for allocations { CUDF_EXPECTS(col1.child().type().id() == cudf::type_id::STRUCT, diff --git a/src/main/cpp/src/map_zip_with_utils.hpp b/src/main/cpp/src/map_zip_with_utils.hpp index 6edc7442b0..148583c6dd 100644 --- a/src/main/cpp/src/map_zip_with_utils.hpp +++ b/src/main/cpp/src/map_zip_with_utils.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2025, NVIDIA CORPORATION. + * Copyright (c) 2025-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -18,6 +18,8 @@ #include #include +#include + namespace spark_rapids_jni { /** @@ -60,7 +62,7 @@ namespace spark_rapids_jni { [[maybe_unused]] std::unique_ptr map_zip( cudf::lists_column_view const& col1, cudf::lists_column_view const& col2, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni \ No newline at end of file diff --git a/src/main/cpp/src/multiply.cu b/src/main/cpp/src/multiply.cu index 85a1d400d4..08e9154c95 100644 --- a/src/main/cpp/src/multiply.cu +++ b/src/main/cpp/src/multiply.cu @@ -181,7 +181,7 @@ std::unique_ptr multiply_impl(cudf::data_type type, RIGHT_ACCESSOR right_accessor, bool check_overflow, bool both_inputs_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto result = @@ -239,7 +239,7 @@ struct dispatch_multiply { std::unique_ptr operator()(cudf::data_type type, cudf::size_type num_rows, bool check_overflow, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { if (left_cv != nullptr && right_cv != nullptr) { @@ -307,7 +307,7 @@ struct dispatch_multiply { std::unique_ptr operator()(cudf::data_type type, cudf::size_type num_rows, bool check_overflow, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { CUDF_FAIL("Unsupported type when multiply."); @@ -320,7 +320,7 @@ std::unique_ptr multiply(cudf::column_view const& left_cv, cudf::column_view const& right_cv, bool is_ansi_mode, bool is_try_mode, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { check_multiply_inputs( @@ -344,7 +344,7 @@ std::unique_ptr multiply(cudf::column_view const& left_cv, cudf::scalar const& right_scalar, bool is_ansi_mode, bool is_try_mode, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { check_multiply_inputs( @@ -368,7 +368,7 @@ std::unique_ptr multiply(cudf::scalar const& left_scalar, cudf::column_view const& right_cv, bool is_ansi_mode, bool is_try_mode, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { check_multiply_inputs(left_scalar.type(), diff --git a/src/main/cpp/src/multiply.hpp b/src/main/cpp/src/multiply.hpp index c5b6158e4e..0b729e19b4 100644 --- a/src/main/cpp/src/multiply.hpp +++ b/src/main/cpp/src/multiply.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2025, NVIDIA CORPORATION. + * Copyright (c) 2025-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -20,6 +20,8 @@ #include #include +#include + namespace spark_rapids_jni { /** @@ -51,7 +53,7 @@ std::unique_ptr multiply( cudf::column_view const& right_input, bool is_ansi_mode, bool is_try_mode, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -84,7 +86,7 @@ std::unique_ptr multiply( cudf::scalar const& right_input, bool is_ansi_mode, bool is_try_mode, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -117,7 +119,7 @@ std::unique_ptr multiply( cudf::column_view const& right_input, bool is_ansi_mode, bool is_try_mode, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/number_converter.cu b/src/main/cpp/src/number_converter.cu index 631543b789..5833aefe07 100644 --- a/src/main/cpp/src/number_converter.cu +++ b/src/main/cpp/src/number_converter.cu @@ -368,7 +368,7 @@ std::unique_ptr convert_impl(cudf::size_type num_rows, STR_ITERATOR input, FROM_BASE_ITERATOR from_base, TO_BASE_ITERATOR to_base, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { static constexpr bool IS_CONST_BASES = cuda::std::is_same_v && @@ -454,7 +454,7 @@ bool is_convert_overflow_impl(cudf::size_type num_rows, STR_ITERATOR input, FROM_BASE_ITERATOR from_base, TO_BASE_ITERATOR to_base, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { static constexpr bool IS_CONST_BASES = cuda::std::is_same_v && @@ -479,7 +479,7 @@ bool is_cv(convert_number_t const& t) { return std::holds_alternative convert(convert_number_t const& input, convert_number_t const& from_base, convert_number_t const& to_base, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { check_types(input, from_base, to_base, stream); @@ -627,7 +627,7 @@ std::unique_ptr convert(convert_number_t const& input, bool is_convert_overflow(convert_number_t const& input, convert_number_t const& from_base, convert_number_t const& to_base, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { check_types(input, from_base, to_base, stream); diff --git a/src/main/cpp/src/number_converter.hpp b/src/main/cpp/src/number_converter.hpp index bbb569590a..4e6b9b36b3 100644 --- a/src/main/cpp/src/number_converter.hpp +++ b/src/main/cpp/src/number_converter.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2025, NVIDIA CORPORATION. + * Copyright (c) 2025-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -20,6 +20,8 @@ #include #include +#include + #include #include @@ -49,7 +51,7 @@ std::unique_ptr convert( convert_number_t const& input, convert_number_t const& from_base, convert_number_t const& to_base, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -68,7 +70,7 @@ bool is_convert_overflow( convert_number_t const& input, convert_number_t const& from_base, convert_number_t const& to_base, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/parse_timestamp_with_format.cu b/src/main/cpp/src/parse_timestamp_with_format.cu index 36f8b417c4..26cfa9adac 100644 --- a/src/main/cpp/src/parse_timestamp_with_format.cu +++ b/src/main/cpp/src/parse_timestamp_with_format.cu @@ -27,10 +27,10 @@ #include #include -#include #include #include +#include #include #include @@ -337,7 +337,7 @@ std::unique_ptr parse_timestamp_strings_with_format( cudf::strings_column_view const& input, std::string const& format, bool legacy, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -353,7 +353,7 @@ std::unique_ptr parse_timestamp_strings_with_format( host_tokens.data(), sizeof(format_token) * host_tokens.size(), cudaMemcpyDefault, - stream.value())); + stream.get())); auto const d_input = cudf::column_device_view::create( input.parent(), stream, cudf::get_current_device_resource_ref()); diff --git a/src/main/cpp/src/parse_uri.cu b/src/main/cpp/src/parse_uri.cu index cbdf93ef07..d391785158 100644 --- a/src/main/cpp/src/parse_uri.cu +++ b/src/main/cpp/src/parse_uri.cu @@ -32,12 +32,12 @@ #include #include -#include #include #include #include #include +#include #include #include @@ -887,7 +887,7 @@ CUDF_KERNEL void parse_uri(column_device_view const in_strings, std::unique_ptr parse_uri(strings_column_view const& input, URI_chunks chunk, std::optional query_match, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { size_type strings_count = input.size(); @@ -920,7 +920,7 @@ std::unique_ptr parse_uri(strings_column_view const& input, // count number of bytes in each string after parsing and store it in offsets_column auto offsets_view = offsets_column->view(); auto offsets_mutable_view = offsets_column->mutable_view(); - parse_uri_char_counter<<>>( + parse_uri_char_counter<<>>( *d_strings, chunk, input.chars_begin(stream), @@ -946,7 +946,7 @@ std::unique_ptr parse_uri(strings_column_view const& input, 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<<>>( + parse_uri<<>>( *d_strings, input.chars_begin(stream), src_offsets.data(), @@ -963,7 +963,7 @@ std::unique_ptr parse_uri(strings_column_view const& input, std::move(null_mask)); } -void validate_input_uris(strings_column_view const& input, rmm::cuda_stream_view stream) +void validate_input_uris(strings_column_view const& input, cuda::stream_ref stream) { if (input.size() == 0) { return; } @@ -1007,7 +1007,7 @@ void validate_input_uris(strings_column_view const& input, rmm::cuda_stream_view std::unique_ptr parse_uri_to_protocol(strings_column_view const& input, bool ansi_mode, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -1017,7 +1017,7 @@ std::unique_ptr parse_uri_to_protocol(strings_column_view const& input, std::unique_ptr parse_uri_to_host(strings_column_view const& input, bool ansi_mode, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -1027,7 +1027,7 @@ std::unique_ptr parse_uri_to_host(strings_column_view const& input, std::unique_ptr parse_uri_to_query(strings_column_view const& input, bool ansi_mode, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -1038,7 +1038,7 @@ std::unique_ptr parse_uri_to_query(strings_column_view const& input, std::unique_ptr parse_uri_to_query(cudf::strings_column_view const& input, std::string const& query_match, bool ansi_mode, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -1054,7 +1054,7 @@ std::unique_ptr parse_uri_to_query(cudf::strings_column_view const std::unique_ptr parse_uri_to_query(cudf::strings_column_view const& input, cudf::strings_column_view const& query_match, bool ansi_mode, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -1066,7 +1066,7 @@ std::unique_ptr parse_uri_to_query(cudf::strings_column_view const std::unique_ptr parse_uri_to_path(strings_column_view const& input, bool ansi_mode, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/parse_uri.hpp b/src/main/cpp/src/parse_uri.hpp index 9262e57ccb..a06392fc38 100644 --- a/src/main/cpp/src/parse_uri.hpp +++ b/src/main/cpp/src/parse_uri.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2025, NVIDIA CORPORATION. + * Copyright (c) 2023-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -20,9 +20,10 @@ #include #include -#include #include +#include + #include namespace spark_rapids_jni { @@ -38,7 +39,7 @@ namespace spark_rapids_jni { std::unique_ptr parse_uri_to_protocol( cudf::strings_column_view const& input, bool ansi_mode = false, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -52,7 +53,7 @@ std::unique_ptr parse_uri_to_protocol( std::unique_ptr parse_uri_to_host( cudf::strings_column_view const& input, bool ansi_mode = false, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -66,7 +67,7 @@ std::unique_ptr parse_uri_to_host( std::unique_ptr parse_uri_to_query( cudf::strings_column_view const& input, bool ansi_mode = false, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -82,7 +83,7 @@ std::unique_ptr parse_uri_to_query( cudf::strings_column_view const& input, std::string const& query_match, bool ansi_mode = false, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -98,7 +99,7 @@ std::unique_ptr parse_uri_to_query( cudf::strings_column_view const& input, cudf::strings_column_view const& query_match, bool ansi_mode = false, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); /** @@ -112,7 +113,7 @@ std::unique_ptr parse_uri_to_query( std::unique_ptr parse_uri_to_path( cudf::strings_column_view const& input, bool ansi_mode = false, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/protobuf/protobuf.cu b/src/main/cpp/src/protobuf/protobuf.cu index 18c850074d..9f8d38b4e8 100644 --- a/src/main/cpp/src/protobuf/protobuf.cu +++ b/src/main/cpp/src/protobuf/protobuf.cu @@ -42,14 +42,14 @@ namespace detail { namespace { void propagate_nulls_to_descendants(cudf::column& col, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); void apply_parent_mask_to_row_aligned_column(cudf::column& col, cudf::bitmask_type const* parent_mask_ptr, cudf::size_type parent_null_count, cudf::size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (parent_null_count == 0) { return; } @@ -78,7 +78,7 @@ void apply_parent_mask_to_row_aligned_column(cudf::column& col, } void propagate_list_nulls_to_descendants(cudf::column& list_col, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (list_col.type().id() != cudf::type_id::LIST || list_col.null_count() == 0) { return; } @@ -118,7 +118,7 @@ void propagate_list_nulls_to_descendants(cudf::column& list_col, } void propagate_struct_nulls_to_descendants(cudf::column& struct_col, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (struct_col.type().id() != cudf::type_id::STRUCT || struct_col.null_count() == 0) { return; } @@ -137,7 +137,7 @@ void propagate_struct_nulls_to_descendants(cudf::column& struct_col, } void propagate_nulls_to_descendants(cudf::column& col, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { switch (col.type().id()) { @@ -152,7 +152,7 @@ void propagate_nulls_to_descendants(cudf::column& col, std::unique_ptr make_null_column_with_schema(protobuf_schema const& schema, int schema_idx, cudf::size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const& field = schema[schema_idx]; @@ -404,7 +404,7 @@ bool protobuf_schema::is_output(int schema_idx) const std::unique_ptr decode_protobuf_to_struct(cudf::column_view const& binary_input, protobuf_decode_context const& context, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { protobuf_schema schema_context{context}; @@ -485,7 +485,7 @@ std::unique_ptr decode_protobuf_to_struct(cudf::column_view const& track_permissive_null_rows ? num_rows : 0, stream, cudf::get_current_device_resource_ref()); if (track_permissive_null_rows) { CUDF_CUDA_TRY( - cudaMemsetAsync(d_row_force_null.data(), 0, num_rows * sizeof(bool), stream.value())); + cudaMemsetAsync(d_row_force_null.data(), 0, num_rows * sizeof(bool), stream.get())); } auto const decode_ctx = protobuf_decode_runtime_context{&d_row_force_null, &d_error}; @@ -627,7 +627,7 @@ std::unique_ptr decode_protobuf_to_struct(cudf::column_view const& auto const batch_input = batched_scalar_input_view{ input, d_locations.data(), num_scalar, d_descs.data(), nf, d_error.data()}; extract_scalar_batched_kernel - <<>>(batch_input); + <<>>(batch_input); } for (int j = 0; j < nf; j++) { @@ -838,7 +838,7 @@ std::unique_ptr decode_protobuf_to_struct(cudf::column_view const& protobuf_error h_error = NONE; CUDF_CUDA_TRY( cudf::detail::memcpy_async(&h_error, d_error.data(), sizeof(protobuf_error), stream)); - stream.synchronize(); + stream.sync(); if (h_error == SCHEMA_TOO_LARGE || h_error == REPEATED_COUNT_MISMATCH) { throw cudf::logic_error(error_message(h_error)); } @@ -891,7 +891,7 @@ std::unique_ptr decode_protobuf_to_struct(cudf::column_view const& std::unique_ptr decode_protobuf_to_struct(cudf::column_view const& binary_input, protobuf_decode_context const& context, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/protobuf/protobuf.hpp b/src/main/cpp/src/protobuf/protobuf.hpp index 2975ebb4ea..968238c918 100644 --- a/src/main/cpp/src/protobuf/protobuf.hpp +++ b/src/main/cpp/src/protobuf/protobuf.hpp @@ -21,9 +21,10 @@ #include #include -#include #include +#include + #include #include #include @@ -98,7 +99,7 @@ void validate_decode_context(protobuf_decode_context const& context); std::unique_ptr decode_protobuf_to_struct(cudf::column_view const& binary_input, protobuf_decode_context const& context, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace spark_rapids_jni::protobuf diff --git a/src/main/cpp/src/protobuf/protobuf_builders.cu b/src/main/cpp/src/protobuf/protobuf_builders.cu index ae932db7ba..95e0d6f27e 100644 --- a/src/main/cpp/src/protobuf/protobuf_builders.cu +++ b/src/main/cpp/src/protobuf/protobuf_builders.cu @@ -35,7 +35,7 @@ namespace spark_rapids_jni::protobuf::detail { field_descriptor_bundle make_field_descriptors(std::vector const& field_indices, protobuf_schema const& schema, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr, std::span output_indices) { @@ -58,7 +58,7 @@ namespace { inline std::pair make_null_mask_from_parent_locations( field_location const* parent_locs, int num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(num_rows >= 0, std::string{__func__} + ": row count must be non-negative"); @@ -106,7 +106,7 @@ inline std::unique_ptr make_list_column_with_parent_nulls( std::unique_ptr offsets_col, std::unique_ptr child_col, field_location const* parent_locs, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto [list_mask, list_null_count] = @@ -122,7 +122,7 @@ std::unique_ptr make_list_column_with_input_nulls( std::unique_ptr offsets_col, std::unique_ptr child_col, cudf::column_view const& binary_input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const input_null_count = binary_input.null_count(); @@ -139,7 +139,7 @@ std::unique_ptr make_list_column_with_input_nulls( std::unique_ptr make_null_column(cudf::data_type dtype, cudf::size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (num_rows == 0) { return cudf::make_empty_column(dtype); } @@ -179,7 +179,7 @@ std::unique_ptr make_null_column(cudf::data_type dtype, } std::unique_ptr make_empty_column_safe(cudf::data_type dtype, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { switch (dtype.id()) { @@ -191,7 +191,7 @@ std::unique_ptr make_empty_column_safe(cudf::data_type dtype, rmm::device_buffer{}, 0); CUDF_CUDA_TRY(cudaMemsetAsync( - offsets_col->mutable_view().data(), 0, sizeof(int32_t), stream.value())); + offsets_col->mutable_view().data(), 0, sizeof(int32_t), stream.get())); auto child_col = std::make_unique( cudf::data_type{cudf::type_id::UINT8}, 0, rmm::device_buffer{}, rmm::device_buffer{}, 0); return cudf::make_lists_column( @@ -209,7 +209,7 @@ std::unique_ptr make_empty_column_safe(cudf::data_type dtype, std::unique_ptr make_null_list_column_with_child( std::unique_ptr child_col, cudf::size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { rmm::device_uvector offsets(num_rows + 1, stream, mr); @@ -221,7 +221,7 @@ std::unique_ptr make_null_list_column_with_child( } std::unique_ptr make_empty_list_column(std::unique_ptr element_col, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto offsets_col = std::make_unique(cudf::data_type{cudf::type_id::INT32}, @@ -229,8 +229,8 @@ std::unique_ptr make_empty_list_column(std::unique_ptrmutable_view().data(), 0, sizeof(int32_t), stream.value())); + CUDF_CUDA_TRY( + cudaMemsetAsync(offsets_col->mutable_view().data(), 0, sizeof(int32_t), stream.get())); return cudf::make_lists_column( 0, std::move(offsets_col), std::move(element_col), 0, rmm::device_buffer{}); } @@ -255,7 +255,7 @@ struct enum_string_lookup_tables { enum_string_lookup_tables make_enum_string_lookup_tables( cudf::detail::host_vector const& valid_enums, std::vector> const& enum_name_bytes, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto d_valid_enums = cudf::detail::make_device_uvector_async( @@ -301,7 +301,7 @@ std::unique_ptr build_enum_string_values_column( rmm::device_uvector& valid, enum_string_lookup_tables const& lookup, int num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { rmm::device_uvector lengths(num_rows, stream, cudf::get_current_device_resource_ref()); @@ -325,7 +325,7 @@ std::unique_ptr build_enum_string_values_column( std::unique_ptr build_enum_string_column(rmm::device_uvector& enum_values, rmm::device_uvector& valid, protobuf_field_decode_request request, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const field = request.context.schema.field(request.schema_idx); @@ -342,7 +342,7 @@ std::unique_ptr build_repeated_enum_string_column( protobuf_input_view input, recursive_decode_context context, repeated_field_work work, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { validate_nonempty_repeated_field_work(work, input.num_rows); @@ -385,7 +385,7 @@ std::unique_ptr build_repeated_string_column( repeated_field_work work, bool is_bytes, rmm::device_uvector& d_error, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { validate_nonempty_repeated_field_work(work, input.num_rows); @@ -399,7 +399,7 @@ std::unique_ptr build_repeated_string_column( repeated_location_provider loc_provider{ input.row_offsets, input.base_offset, work.occurrences->data()}; extract_lengths_kernel - <<>>(loc_provider, total_count, str_lengths.data()); + <<>>(loc_provider, total_count, str_lengths.data()); auto [str_offsets_col, total_chars] = cudf::strings::detail::make_offsets_child_column( str_lengths.begin(), str_lengths.end(), stream, mr); @@ -435,7 +435,7 @@ std::unique_ptr build_repeated_string_column( size_t temp_storage_bytes = 0; cub::DeviceMemcpy::Batched( - nullptr, temp_storage_bytes, src_iter, dst_iter, size_iter, total_count, stream.value()); + nullptr, temp_storage_bytes, src_iter, dst_iter, size_iter, total_count, stream.get()); rmm::device_buffer temp_storage(temp_storage_bytes, stream, scratch_mr); cub::DeviceMemcpy::Batched(temp_storage.data(), temp_storage_bytes, @@ -443,7 +443,7 @@ std::unique_ptr build_repeated_string_column( dst_iter, size_iter, total_count, - stream.value()); + stream.get()); } std::unique_ptr child_col; @@ -482,7 +482,7 @@ std::unique_ptr build_nested_struct_column( std::vector const& child_field_indices, recursive_decode_context context, int depth, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const& schema = context.schema; @@ -628,7 +628,7 @@ std::unique_ptr build_repeated_child_list_column(protobuf_input_vi nested_parent_view parent, recursive_decode_context context, repeated_field_work work, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const& schema = context.schema; diff --git a/src/main/cpp/src/protobuf/protobuf_device_helpers.cuh b/src/main/cpp/src/protobuf/protobuf_device_helpers.cuh index 6ad4787d77..6b8591b332 100644 --- a/src/main/cpp/src/protobuf/protobuf_device_helpers.cuh +++ b/src/main/cpp/src/protobuf/protobuf_device_helpers.cuh @@ -20,11 +20,10 @@ #include -#include - #include #include #include +#include #include #include @@ -80,7 +79,7 @@ __device__ __forceinline__ void write_varint_value(T* dst, uint64_t val) void set_error_once_async(protobuf_error* error_flag, protobuf_error error, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); __device__ inline int get_wire_type_size(proto_wire_type wt, uint8_t const* cur, uint8_t const* end) { diff --git a/src/main/cpp/src/protobuf/protobuf_host_helpers.hpp b/src/main/cpp/src/protobuf/protobuf_host_helpers.hpp index 622d914cf0..9008c87c30 100644 --- a/src/main/cpp/src/protobuf/protobuf_host_helpers.hpp +++ b/src/main/cpp/src/protobuf/protobuf_host_helpers.hpp @@ -24,12 +24,12 @@ #include #include -#include #include #include #include #include +#include #include #include #include @@ -93,7 +93,7 @@ struct field_descriptor_bundle { field_descriptor_bundle make_field_descriptors(std::vector const& field_indices, protobuf_schema const& schema, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr, std::span output_indices = {}); @@ -175,7 +175,7 @@ inline list_offsets_from_counts_result make_list_offsets_from_counts( CountIterator counts_begin, int num_rows, char const* count_context, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref output_mr, rmm::device_async_resource_ref scratch_mr) { @@ -223,7 +223,7 @@ inline repeated_field_work_bundle make_repeated_field_work_bundle( int num_rows, protobuf_schema const& schema, char const* count_context, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref output_mr, rmm::device_async_resource_ref scratch_mr) { @@ -262,7 +262,7 @@ inline repeated_field_work_bundle make_repeated_field_work_bundle( inline rmm::device_uvector make_top_row_indices( rmm::device_uvector const& occurrences, int32_t const* parent_top_row_indices, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { rmm::device_uvector result(occurrences.size(), stream, mr); @@ -291,7 +291,7 @@ inline rmm::device_uvector make_top_row_indices( template inline cudf::detail::host_vector build_lookup_table(FieldNumberFn get_field_number, int num_entries, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (num_entries == 0) { return cudf::detail::make_pinned_vector_async(0, stream); } @@ -313,7 +313,7 @@ inline cudf::detail::host_vector build_lookup_table(FieldNumberFn get_field template inline cudf::detail::host_vector build_field_lookup_table(FieldDesc const* descs, int num_fields, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { return build_lookup_table([&](int i) { return descs[i].field_number; }, num_fields, stream); } @@ -334,7 +334,7 @@ struct field_occurrence_scan_bundle { inline field_occurrence_scan_bundle make_field_occurrence_scan_bundle( cudf::detail::host_vector const& host_descriptors, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto descriptors = cudf::detail::make_device_uvector_async(host_descriptors, stream, mr); @@ -371,11 +371,11 @@ inline std::vector const& find_child_field_indices(protobuf_schema const& s // Forward declarations needed by make_empty_struct_column_with_schema std::unique_ptr make_empty_column_safe(cudf::data_type dtype, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr make_empty_list_column(std::unique_ptr element_col, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); // Forward declaration for the mutual recursion with make_empty_struct_column_from_children. @@ -383,7 +383,7 @@ template std::unique_ptr make_empty_struct_column_with_schema( SchemaT const& schema, int parent_idx, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); // Build an empty (0-row) STRUCT column from an explicit child-index list, recursing into @@ -393,7 +393,7 @@ template std::unique_ptr make_empty_struct_column_from_children( SchemaT const& schema, std::vector const& child_indices, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { std::vector> children; @@ -419,10 +419,7 @@ std::unique_ptr make_empty_struct_column_from_children( template std::unique_ptr make_empty_struct_column_with_schema( - SchemaT const& schema, - int parent_idx, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) + SchemaT const& schema, int parent_idx, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const& child_indices = find_child_field_indices(schema, parent_idx); return make_empty_struct_column_from_children(schema, child_indices, stream, mr); @@ -432,26 +429,26 @@ void maybe_check_required_fields(required_field_input_view input, std::vector const& field_indices, std::vector const& schema, protobuf_decode_runtime_context decode_ctx, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); void propagate_invalid_enum_flags_to_rows(rmm::device_uvector const& item_invalid, protobuf_decode_runtime_context decode_ctx, protobuf_value_domain_view value_domain, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); void validate_enum_and_propagate_rows(rmm::device_uvector const& values, rmm::device_uvector& valid, enum_domain_device_view enum_domain, protobuf_decode_runtime_context decode_ctx, protobuf_value_domain_view value_domain, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); void validate_enum_and_propagate_rows(rmm::device_uvector const& values, rmm::device_uvector& valid, cudf::detail::host_vector const& valid_enums, protobuf_decode_runtime_context decode_ctx, protobuf_value_domain_view value_domain, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); // ============================================================================ // Forward declarations of builder/utility functions @@ -459,7 +456,7 @@ void validate_enum_and_propagate_rows(rmm::device_uvector const& values std::unique_ptr make_null_column(cudf::data_type dtype, cudf::size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); // Schema-aware all-null builder: recurses into STRUCT children and wraps repeated fields @@ -468,19 +465,19 @@ std::unique_ptr make_null_column(cudf::data_type dtype, std::unique_ptr make_null_column_with_schema(protobuf_schema const& schema, int schema_idx, cudf::size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr make_null_list_column_with_child( std::unique_ptr child_col, cudf::size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr build_enum_string_column(rmm::device_uvector& enum_values, rmm::device_uvector& valid, protobuf_field_decode_request request, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); // Wrap offsets + child into a LIST column, propagating the input's null mask. Note: when @@ -491,7 +488,7 @@ std::unique_ptr make_list_column_with_input_nulls( std::unique_ptr offsets_col, std::unique_ptr child_col, cudf::column_view const& binary_input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr build_repeated_enum_string_column( @@ -499,7 +496,7 @@ std::unique_ptr build_repeated_enum_string_column( protobuf_input_view input, recursive_decode_context context, repeated_field_work work, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr build_repeated_string_column( @@ -508,7 +505,7 @@ std::unique_ptr build_repeated_string_column( repeated_field_work work, bool is_bytes, rmm::device_uvector& d_error, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr build_nested_struct_column( @@ -517,14 +514,14 @@ std::unique_ptr build_nested_struct_column( std::vector const& child_field_indices, recursive_decode_context context, int depth, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr build_repeated_child_list_column(protobuf_input_view input, nested_parent_view parent, recursive_decode_context context, repeated_field_work work, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace spark_rapids_jni::protobuf::detail diff --git a/src/main/cpp/src/protobuf/protobuf_kernels.cu b/src/main/cpp/src/protobuf/protobuf_kernels.cu index 8b870d4bfe..93adc10d14 100644 --- a/src/main/cpp/src/protobuf/protobuf_kernels.cu +++ b/src/main/cpp/src/protobuf/protobuf_kernels.cu @@ -737,11 +737,9 @@ CUDF_KERNEL void copy_enum_string_chars_kernel(enum_value_device_view input, // Host wrapper functions — callable from other translation units // ============================================================================ -void set_error_once_async(protobuf_error* error_flag, - protobuf_error error, - rmm::cuda_stream_view stream) +void set_error_once_async(protobuf_error* error_flag, protobuf_error error, cuda::stream_ref stream) { - set_error_if_unset_kernel<<<1, 1, 0, stream.value()>>>(error_flag, error); + set_error_if_unset_kernel<<<1, 1, 0, stream.get()>>>(error_flag, error); CUDF_CUDA_TRY(cudaPeekAtLastError()); } @@ -749,12 +747,12 @@ void launch_scan_all_fields(cudf::column_device_view const& d_in, field_scan_view fields, protobuf_error* error_flag, bool* row_has_invalid_data, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const num_rows = d_in.size(); if (num_rows == 0) return; auto const blocks = static_cast((num_rows + THREADS_PER_BLOCK - 1u) / THREADS_PER_BLOCK); - scan_all_fields_kernel<<>>( + scan_all_fields_kernel<<>>( d_in, fields, error_flag, row_has_invalid_data); } @@ -762,24 +760,24 @@ void launch_count_repeated_fields(cudf::column_device_view const& d_in, field_scan_view fields, protobuf_error* error_flag, bool* row_has_invalid_data, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const num_rows = d_in.size(); if (num_rows == 0) return; auto const blocks = static_cast((num_rows + THREADS_PER_BLOCK - 1u) / THREADS_PER_BLOCK); - count_repeated_fields_kernel<<>>( + count_repeated_fields_kernel<<>>( d_in, fields, error_flag, row_has_invalid_data); } void launch_scan_all_field_occurrences(cudf::column_device_view const& d_in, field_occurrence_scan_view fields, protobuf_error* error_flag, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const num_rows = d_in.size(); if (num_rows == 0) return; auto const blocks = static_cast((num_rows + THREADS_PER_BLOCK - 1u) / THREADS_PER_BLOCK); - scan_all_field_occurrences_kernel<<>>( + scan_all_field_occurrences_kernel<<>>( d_in, fields, error_flag); } @@ -788,11 +786,11 @@ void launch_extract_strided_locations(field_location const* nested_locations, int num_fields, field_location* parent_locs, int num_rows, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (num_rows == 0) return; auto const blocks = static_cast((num_rows + THREADS_PER_BLOCK - 1u) / THREADS_PER_BLOCK); - extract_strided_locations_kernel<<>>( + extract_strided_locations_kernel<<>>( nested_locations, field_idx, num_fields, parent_locs, num_rows); } @@ -801,12 +799,12 @@ void launch_scan_nested_message_fields(protobuf_input_view input, field_scan_view fields, protobuf_error* error_flag, bool* row_has_invalid_data, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (input.num_rows == 0) return; auto const blocks = static_cast((input.num_rows + THREADS_PER_BLOCK - 1u) / THREADS_PER_BLOCK); - scan_nested_message_fields_kernel<<>>( + scan_nested_message_fields_kernel<<>>( input, parent, fields, error_flag, row_has_invalid_data); } @@ -814,12 +812,12 @@ void launch_scan_all_field_occurrences_in_nested(protobuf_input_view input, nested_parent_view parent, field_occurrence_scan_view fields, protobuf_error* error_flag, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (input.num_rows == 0) return; auto const blocks = static_cast((input.num_rows + THREADS_PER_BLOCK - 1u) / THREADS_PER_BLOCK); - scan_all_field_occurrences_in_nested_kernel<<>>( + scan_all_field_occurrences_in_nested_kernel<<>>( input, parent, fields, error_flag); } @@ -830,33 +828,33 @@ void launch_compute_grandchild_parent_locations(field_location const* parent_loc field_location* gc_parent_locs, int num_rows, protobuf_error* error_flag, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (num_rows == 0) return; auto const blocks = static_cast((num_rows + THREADS_PER_BLOCK - 1u) / THREADS_PER_BLOCK); - compute_grandchild_parent_locations_kernel<<>>( + compute_grandchild_parent_locations_kernel<<>>( parent_locs, child_locs, child_idx, num_child_fields, gc_parent_locs, num_rows, error_flag); } void launch_validate_enum_values(enum_value_device_view input, bool* row_has_invalid_enum, enum_domain_device_view domain, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (input.size == 0) return; auto const blocks = static_cast((input.size + THREADS_PER_BLOCK - 1u) / THREADS_PER_BLOCK); - validate_enum_values_kernel<<>>( + validate_enum_values_kernel<<>>( input, row_has_invalid_enum, domain); } void launch_compute_enum_string_lengths(enum_value_device_view input, enum_string_lookup_device_view lookup, int32_t* lengths, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (input.size == 0) return; auto const blocks = static_cast((input.size + THREADS_PER_BLOCK - 1u) / THREADS_PER_BLOCK); - compute_enum_string_lengths_kernel<<>>( + compute_enum_string_lengths_kernel<<>>( input, lookup, lengths); } @@ -864,11 +862,11 @@ void launch_copy_enum_string_chars(enum_value_device_view input, enum_string_lookup_device_view lookup, int32_t const* output_offsets, char* out_chars, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (input.size == 0) return; auto const blocks = static_cast((input.size + THREADS_PER_BLOCK - 1u) / THREADS_PER_BLOCK); - copy_enum_string_chars_kernel<<>>( + copy_enum_string_chars_kernel<<>>( input, lookup, output_offsets, out_chars); } @@ -876,7 +874,7 @@ void maybe_check_required_fields(required_field_input_view input, std::vector const& field_indices, std::vector const& schema, protobuf_decode_runtime_context decode_ctx, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (input.values.size == 0 || field_indices.empty()) { return; } @@ -895,7 +893,7 @@ void maybe_check_required_fields(required_field_input_view input, auto const blocks = static_cast((input.values.size + THREADS_PER_BLOCK - 1u) / THREADS_PER_BLOCK); - check_required_fields_kernel<<>>( + check_required_fields_kernel<<>>( input, d_is_required.data(), static_cast(field_indices.size()), @@ -906,7 +904,7 @@ void maybe_check_required_fields(required_field_input_view input, void propagate_invalid_enum_flags_to_rows(rmm::device_uvector const& item_invalid, protobuf_decode_runtime_context decode_ctx, protobuf_value_domain_view value_domain, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto& row_invalid = *decode_ctx.row_force_null; auto const num_items = value_domain.size; @@ -952,7 +950,7 @@ void validate_enum_and_propagate_rows(rmm::device_uvector const& values enum_domain_device_view enum_domain, protobuf_decode_runtime_context decode_ctx, protobuf_value_domain_view value_domain, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (value_domain.size == 0 || enum_domain.size == 0) return; @@ -971,7 +969,7 @@ void validate_enum_and_propagate_rows(rmm::device_uvector const& values cudf::detail::host_vector const& valid_enums, protobuf_decode_runtime_context decode_ctx, protobuf_value_domain_view value_domain, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (value_domain.size == 0 || valid_enums.empty()) return; diff --git a/src/main/cpp/src/protobuf/protobuf_kernels.cuh b/src/main/cpp/src/protobuf/protobuf_kernels.cuh index 89ce209963..d713ccea36 100644 --- a/src/main/cpp/src/protobuf/protobuf_kernels.cuh +++ b/src/main/cpp/src/protobuf/protobuf_kernels.cuh @@ -39,6 +39,7 @@ #include #include #include +#include #include #include #include @@ -411,32 +412,32 @@ void launch_count_repeated_fields(cudf::column_device_view const& d_in, field_scan_view fields, protobuf_error* error_flag, bool* row_has_invalid_data, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); void launch_scan_all_field_occurrences(cudf::column_device_view const& d_in, field_occurrence_scan_view fields, protobuf_error* error_flag, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); void launch_extract_strided_locations(field_location const* nested_locations, int field_idx, int num_fields, field_location* parent_locs, int num_rows, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); void launch_scan_nested_message_fields(protobuf_input_view input, nested_parent_view parent, field_scan_view fields, protobuf_error* error_flag, bool* row_has_invalid_data, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); void launch_scan_all_field_occurrences_in_nested(protobuf_input_view input, nested_parent_view parent, field_occurrence_scan_view fields, protobuf_error* error_flag, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); void launch_compute_grandchild_parent_locations(field_location const* parent_locs, field_location const* child_locs, @@ -445,7 +446,7 @@ void launch_compute_grandchild_parent_locations(field_location const* parent_loc field_location* gc_parent_locs, int num_rows, protobuf_error* error_flag, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); // ============================================================================ // Host-side template helpers that launch CUDA kernels @@ -456,7 +457,7 @@ template inline std::pair make_null_mask_from_valid( rmm::device_uvector const& valid, cudf::size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(num_rows >= 0, "num_rows must be non-negative"); @@ -477,13 +478,13 @@ inline void extract_scalar_into_buffers(uint8_t const* message_data, proto_encoding encoding, scalar_decode_options options, scalar_value_output output, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto constexpr threads = THREADS_PER_BLOCK; auto const blocks = static_cast((num_rows + threads - 1u) / threads); dispatch_scalar_decoder(get_scalar_decode_kind(encoding), [&]() { - extract_scalar_kernel<<>>( - message_data, loc_provider, num_rows, output, options); + extract_scalar_kernel + <<>>(message_data, loc_provider, num_rows, output, options); }); } @@ -508,7 +509,7 @@ std::unique_ptr extract_and_build_scalar_field_column( LocationProvider const& loc_provider, int num_rows, protobuf_decode_runtime_context decode_ctx, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (num_rows == 0) { return cudf::make_empty_column(field.output_type); } @@ -537,7 +538,7 @@ inline std::unique_ptr extract_and_build_string_or_bytes_column( bool has_default, cudf::detail::host_vector const& default_bytes, rmm::device_uvector& d_error, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { int32_t def_len = has_default ? static_cast(default_bytes.size()) : 0; @@ -550,7 +551,7 @@ inline std::unique_ptr extract_and_build_string_or_bytes_column( rmm::device_uvector lengths(num_rows, stream, mr); auto const threads = THREADS_PER_BLOCK; auto const blocks = static_cast((num_rows + threads - 1u) / threads); - extract_lengths_kernel<<>>( + extract_lengths_kernel<<>>( loc_provider, num_rows, lengths.data(), has_default, def_len); auto [offsets_col, total_size] = @@ -592,7 +593,7 @@ inline std::unique_ptr extract_and_build_string_or_bytes_column( size_t temp_storage_bytes = 0; cub::DeviceMemcpy::Batched( - nullptr, temp_storage_bytes, src_iter, dst_iter, size_iter, num_rows, stream.value()); + nullptr, temp_storage_bytes, src_iter, dst_iter, size_iter, num_rows, stream.get()); rmm::device_buffer temp_storage(temp_storage_bytes, stream, mr); cub::DeviceMemcpy::Batched(temp_storage.data(), temp_storage_bytes, @@ -600,7 +601,7 @@ inline std::unique_ptr extract_and_build_string_or_bytes_column( dst_iter, size_iter, num_rows, - stream.value()); + stream.get()); } if (num_rows == 0) { @@ -639,7 +640,7 @@ inline std::unique_ptr extract_and_build_string_or_bytes_column( template inline std::unique_ptr extract_typed_column(protobuf_field_decode_request request, LocationProvider const& loc_provider, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const field = request.context.schema.field(request.schema_idx); @@ -698,7 +699,7 @@ inline std::unique_ptr build_protobuf_field_values_column( LocationProvider const& loc_provider, ValidityFn validity_fn, TopRowIndexProvider get_top_row_indices, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const message_data = request.message_data; @@ -778,7 +779,7 @@ inline std::unique_ptr build_repeated_scalar_column( protobuf_field_meta_view field, repeated_field_work work, rmm::device_uvector& d_error, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { validate_nonempty_repeated_field_work(work, input.num_rows); @@ -811,22 +812,22 @@ void launch_scan_all_fields(cudf::column_device_view const& d_in, field_scan_view fields, protobuf_error* error_flag, bool* row_has_invalid_data, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); void launch_validate_enum_values(enum_value_device_view input, bool* row_has_invalid_enum, enum_domain_device_view domain, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); void launch_compute_enum_string_lengths(enum_value_device_view input, enum_string_lookup_device_view lookup, int32_t* lengths, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); void launch_copy_enum_string_chars(enum_value_device_view input, enum_string_lookup_device_view lookup, int32_t const* output_offsets, char* out_chars, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); } // namespace spark_rapids_jni::protobuf::detail diff --git a/src/main/cpp/src/regex_rewrite_utils.cu b/src/main/cpp/src/regex_rewrite_utils.cu index bfc42dcd16..c0913995b3 100644 --- a/src/main/cpp/src/regex_rewrite_utils.cu +++ b/src/main/cpp/src/regex_rewrite_utils.cu @@ -26,9 +26,9 @@ #include #include -#include #include +#include #include #include @@ -69,7 +69,7 @@ std::unique_ptr find_literal_range_pattern(cudf::strings_column_vi int const range_len, int const start, int const end, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const strings_count = strings.size(); @@ -126,7 +126,7 @@ std::unique_ptr literal_range_pattern(cudf::strings_column_view co int const range_len, int const start, int const end, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/regex_rewrite_utils.hpp b/src/main/cpp/src/regex_rewrite_utils.hpp index ed3874e2bf..8b83ec7d10 100644 --- a/src/main/cpp/src/regex_rewrite_utils.hpp +++ b/src/main/cpp/src/regex_rewrite_utils.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -20,6 +20,8 @@ #include #include +#include + namespace spark_rapids_jni { /** * @brief Check if input string contains regex pattern `literal[start-end]{len,}`, which means @@ -40,6 +42,6 @@ std::unique_ptr literal_range_pattern( int const len, int const start, int const end, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, + cuda::stream_ref stream = rmm::cuda_stream_default, rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/reverse_strings.cu b/src/main/cpp/src/reverse_strings.cu index 25c46304bf..2a7ac7a97a 100644 --- a/src/main/cpp/src/reverse_strings.cu +++ b/src/main/cpp/src/reverse_strings.cu @@ -25,11 +25,11 @@ #include #include -#include #include #include #include +#include #include #include @@ -94,7 +94,7 @@ struct reverse_characters_fn { } // namespace std::unique_ptr reverse_strings(cudf::strings_column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.is_empty()) { return cudf::make_empty_column(cudf::type_id::STRING); } @@ -118,7 +118,7 @@ std::unique_ptr reverse_strings(cudf::strings_column_view const& i } // namespace detail std::unique_ptr reverse_strings(cudf::strings_column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/reverse_strings.hpp b/src/main/cpp/src/reverse_strings.hpp index f35e94ea85..acd6452e43 100644 --- a/src/main/cpp/src/reverse_strings.hpp +++ b/src/main/cpp/src/reverse_strings.hpp @@ -20,6 +20,8 @@ #include #include +#include + #include namespace spark_rapids_jni { @@ -42,7 +44,7 @@ namespace spark_rapids_jni { */ std::unique_ptr reverse_strings( cudf::strings_column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/round_float.cu b/src/main/cpp/src/round_float.cu index 91e8de92d6..51e9251b79 100644 --- a/src/main/cpp/src/round_float.cu +++ b/src/main/cpp/src/round_float.cu @@ -27,10 +27,10 @@ #include #include -#include #include #include +#include #include #include #include @@ -100,7 +100,7 @@ struct half_even_negative { template typename RoundFunctor> std::unique_ptr round_with(cudf::column_view const& input, int32_t decimal_places, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(std::is_floating_point_v) { @@ -139,7 +139,7 @@ struct round_type_dispatcher { std::unique_ptr operator()(cudf::column_view const& input, int32_t decimal_places, cudf::rounding_method method, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(std::is_floating_point_v) { @@ -162,7 +162,7 @@ struct round_type_dispatcher { std::unique_ptr round(cudf::column_view const& input, int32_t decimal_places, cudf::rounding_method method, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -260,7 +260,7 @@ template cudf::size_type find_first_overflow_for_integral_type(cudf::column_view const& input, int32_t decimal_places, cudf::rounding_method method, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { static_assert(std::is_integral_v, "T must be an integral type"); @@ -297,7 +297,7 @@ struct find_overflow_dispatcher { cudf::size_type operator()(cudf::column_view const& input, int32_t decimal_places, cudf::rounding_method method, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { if constexpr (std::is_integral_v) { return find_first_overflow_for_integral_type(input, decimal_places, method, stream); @@ -313,7 +313,7 @@ std::unique_ptr round(cudf::column_view const& input, int32_t decimal_places, cudf::rounding_method method, bool is_ansi_mode, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/round_float.hpp b/src/main/cpp/src/round_float.hpp index 67b9d67eff..ccefb0ed59 100644 --- a/src/main/cpp/src/round_float.hpp +++ b/src/main/cpp/src/round_float.hpp @@ -23,9 +23,10 @@ #include #include -#include #include +#include + namespace spark_rapids_jni { /** @@ -66,7 +67,7 @@ std::unique_ptr round( cudf::column_view const& input, int32_t decimal_places = 0, cudf::rounding_method method = cudf::rounding_method::HALF_UP, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -101,7 +102,7 @@ std::unique_ptr round( int32_t decimal_places, cudf::rounding_method method, bool is_ansi_mode, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/row_conversion.cu b/src/main/cpp/src/row_conversion.cu index 7703044c26..f252809345 100644 --- a/src/main/cpp/src/row_conversion.cu +++ b/src/main/cpp/src/row_conversion.cu @@ -36,7 +36,6 @@ #include #include -#include #include #include #include @@ -48,6 +47,7 @@ #include #include #include +#include #include #include @@ -219,7 +219,7 @@ struct batch_data { std::pair, rmm::device_uvector> build_string_row_offsets(table_view const& tbl, size_type fixed_width_and_validity_size, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const num_rows = tbl.num_rows(); rmm::device_uvector d_row_sizes(num_rows, stream); @@ -1269,7 +1269,7 @@ static std::unique_ptr fixed_width_convert_to_rows( rmm::device_uvector& input_nm, scalar const& zero, scalar const& scalar_size_per_row, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { int64_t const total_allocation = size_per_row * num_rows; @@ -1292,7 +1292,7 @@ static std::unique_ptr fixed_width_convert_to_rows( int shared_size = detail::calc_fixed_width_kernel_dims(num_columns, num_rows, size_per_row, blocks, threads); - copy_to_rows_fixed_width_optimized<<>>( + copy_to_rows_fixed_width_optimized<<>>( start_row, num_rows, num_columns, @@ -1507,7 +1507,7 @@ template batch_data build_batches(size_type num_rows, RowSize row_sizes, bool all_fixed_width, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const total_size = @@ -1601,7 +1601,7 @@ batch_data build_batches(size_type num_rows, output_batch_row_offsets.data(), num_rows_in_batch * sizeof(size_type), cudaMemcpyDefault, - stream.value())); + stream.get())); } batch_row_boundaries.push_back(row_end); @@ -1627,7 +1627,7 @@ batch_data build_batches(size_type num_rows, */ int compute_tile_counts(device_span const& batch_row_boundaries, int desired_tile_height, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { size_type const num_batches = batch_row_boundaries.size() - 1; device_uvector num_tiles(num_batches, stream); @@ -1668,7 +1668,7 @@ size_type build_tiles( int column_end, int desired_tile_height, int total_number_of_rows, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { size_type const num_batches = batch_row_boundaries.size() - 1; device_uvector num_tiles(num_batches, stream); @@ -1831,7 +1831,7 @@ std::vector> convert_to_rows( offsetFunctor offset_functor, column_info_s const& column_info, std::optional> variable_width_offsets, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { int device_id; @@ -1948,7 +1948,7 @@ std::vector> convert_to_rows( // blast through the entire table and convert it detail::copy_to_rows - <<>>( + <<>>( num_rows, tbl.num_columns(), shmem_limit_per_tile, @@ -1962,7 +1962,7 @@ std::vector> convert_to_rows( // note that validity gets the entire table and not the fixed-width portion detail::copy_validity_to_rows - <<>>( + <<>>( num_rows, tbl.num_columns(), shmem_limit_per_tile, @@ -2005,7 +2005,7 @@ std::vector> convert_to_rows( // batch_num_rows) so the inner loop guard `row < num_rows` works for every batch, // including batches whose start lies past per-batch row_count. detail::copy_strings_to_rows - <<>>( + <<>>( batch_row_offset + batch_num_rows, variable_width_table.num_columns(), dev_variable_input_data.data(), @@ -2019,7 +2019,7 @@ std::vector> convert_to_rows( // Drain the async H2D uploads above before variable_width_input_data goes out of scope: // CUDA 13+ may read the host source only when the stream executes the copy. - stream.synchronize(); + stream.sync(); } // split up the output buffer into multiple buffers based on row batch sizes and create list of @@ -2054,7 +2054,7 @@ std::vector> convert_to_rows( // Drain the async H2D uploads above before input_data, input_nm, output_data and // validity_tile_infos go out of scope: CUDA 13+ may read the host source only when the stream // executes the copy. - stream.synchronize(); + stream.sync(); return ret; } @@ -2121,7 +2121,7 @@ inline void check_supported_schema(std::vector const& schema, bool fi * @return vector of list columns containing byte columns of the JCUDF row data */ std::vector> convert_to_rows(table_view const& tbl, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -2184,7 +2184,7 @@ std::vector> convert_to_rows(table_view const& tbl, } std::vector> convert_to_rows_fixed_width_optimized( - table_view const& tbl, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + table_view const& tbl, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); // check_supported_columns rejects non-fixed-width columns up front, so the body below does not @@ -2225,11 +2225,11 @@ std::vector> convert_to_rows_fixed_width_optimized( auto dev_input_nm = make_device_uvector_async(input_nm, stream, mr); using ScalarType = scalar_type_t; - auto zero = make_numeric_scalar(data_type(type_id::INT32), stream.value()); + auto zero = make_numeric_scalar(data_type(type_id::INT32), stream.get()); zero->set_valid_async(true, stream); static_cast(zero.get())->set_value(0, stream); - auto step = make_numeric_scalar(data_type(type_id::INT32), stream.value()); + auto step = make_numeric_scalar(data_type(type_id::INT32), stream.get()); step->set_valid_async(true, stream); static_cast(step.get())->set_value(static_cast(size_per_row), stream); @@ -2253,7 +2253,7 @@ std::vector> convert_to_rows_fixed_width_optimized( // Drain the async H2D uploads above before column_start, column_size, input_data and input_nm // go out of scope: CUDA 13+ may read the host source only when the stream executes the copy. - stream.synchronize(); + stream.sync(); return ret; } @@ -2262,7 +2262,7 @@ namespace { /// @brief Calculates and sets null counts for specified columns void fixup_null_counts(std::vector>& output_columns, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { for (auto& col : output_columns) { col->set_null_count(cudf::null_count(col->view().null_mask(), 0, col->size(), stream)); @@ -2282,7 +2282,7 @@ void fixup_null_counts(std::vector>& output_columns, */ std::unique_ptr convert_from_rows(lists_column_view const& input, std::vector const& schema, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -2353,7 +2353,7 @@ std::unique_ptr
convert_from_rows(lists_column_view const& input, auto make_col = [&output_data, &output_nm](data_type type, size_type num_rows, bool include_nm, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto column = make_fixed_width_column(type, @@ -2458,7 +2458,7 @@ std::unique_ptr
convert_from_rows(lists_column_view const& input, detail::fixed_width_row_offset_functor offset_functor(size_per_row); detail::copy_from_rows - <<>>( + <<>>( num_rows, num_columns, shmem_limit_per_tile, @@ -2471,7 +2471,7 @@ std::unique_ptr
convert_from_rows(lists_column_view const& input, child.data()); detail::copy_validity_from_rows - <<>>( + <<>>( num_rows, num_columns, shmem_limit_per_tile, @@ -2485,7 +2485,7 @@ std::unique_ptr
convert_from_rows(lists_column_view const& input, } else { detail::string_row_offset_functor offset_functor(device_span{input.offsets()}); detail::copy_from_rows - <<>>( + <<>>( num_rows, num_columns, shmem_limit_per_tile, @@ -2498,7 +2498,7 @@ std::unique_ptr
convert_from_rows(lists_column_view const& input, child.data()); detail::copy_validity_from_rows - <<>>( + <<>>( num_rows, num_columns, shmem_limit_per_tile, @@ -2545,7 +2545,7 @@ std::unique_ptr
convert_from_rows(lists_column_view const& input, MAX_STRING_BLOCKS)); detail::copy_strings_from_rows - <<>>( + <<>>( offset_functor, dev_string_row_offsets.data(), dev_string_lengths.data(), @@ -2575,7 +2575,7 @@ std::unique_ptr
convert_from_rows(lists_column_view const& input, // Drain the async H2D uploads above before string_col_offset_ptrs and string_data_col_ptrs // go out of scope: CUDA 13+ may read the host source only when the stream executes the copy. - stream.synchronize(); + stream.sync(); } // Set null counts, because output_columns are modified via mutable-view, @@ -2586,14 +2586,14 @@ std::unique_ptr
convert_from_rows(lists_column_view const& input, // Explicitly drain async H2D uploads before the host staging vectors go out of scope // (CUDA 13+ may read the host source only at stream-execution time). Not left to the // incidental sync in fixup_null_counts, which vanishes if null counts move into the kernel. - stream.synchronize(); + stream.sync(); return std::make_unique
(std::move(output_columns)); } std::unique_ptr
convert_from_rows_fixed_width_optimized(lists_column_view const& input, std::vector const& schema, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -2646,7 +2646,7 @@ std::unique_ptr
convert_from_rows_fixed_width_optimized(lists_column_view int shared_size = detail::calc_fixed_width_kernel_dims(num_columns, num_rows, size_per_row, blocks, threads); - detail::copy_from_rows_fixed_width_optimized<<>>( + detail::copy_from_rows_fixed_width_optimized<<>>( num_rows, num_columns, size_per_row, @@ -2664,7 +2664,7 @@ std::unique_ptr
convert_from_rows_fixed_width_optimized(lists_column_view // Explicitly drain async H2D uploads before the host staging vectors go out of scope // (CUDA 13+ may read the host source only at stream-execution time). Not left to the // incidental sync in fixup_null_counts, which vanishes if null counts move into the kernel. - stream.synchronize(); + stream.sync(); return std::make_unique
(std::move(output_columns)); } diff --git a/src/main/cpp/src/row_conversion.hpp b/src/main/cpp/src/row_conversion.hpp index 134a0cef9a..1602968578 100644 --- a/src/main/cpp/src/row_conversion.hpp +++ b/src/main/cpp/src/row_conversion.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2024, NVIDIA CORPORATION. + * Copyright (c) 2022-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -20,9 +20,10 @@ #include #include -#include #include +#include + #include namespace spark_rapids_jni { @@ -30,25 +31,25 @@ namespace spark_rapids_jni { std::vector> convert_to_rows_fixed_width_optimized( cudf::table_view const& tbl, // TODO need something for validity - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); std::vector> convert_to_rows( cudf::table_view const& tbl, // TODO need something for validity - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); std::unique_ptr convert_from_rows_fixed_width_optimized( cudf::lists_column_view const& input, std::vector const& schema, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); std::unique_ptr convert_from_rows( cudf::lists_column_view const& input, std::vector const& schema, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/shuffle_assemble.cu b/src/main/cpp/src/shuffle_assemble.cu index ec8cda0a70..a6ad5199f8 100644 --- a/src/main/cpp/src/shuffle_assemble.cu +++ b/src/main/cpp/src/shuffle_assemble.cu @@ -35,7 +35,6 @@ #include #include -#include #include #include @@ -43,6 +42,7 @@ #include #include #include +#include #include #include #include @@ -176,7 +176,7 @@ int compute_offset_column_info_traverse(cudf::host_span compute_offset_column_info( shuffle_split_metadata const& metadata, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { std::vector offset_info; @@ -295,7 +295,7 @@ std::tuple, assemble_build_column_info(shuffle_split_metadata const& h_global_metadata, cudf::device_span partitions, cudf::device_span partition_offsets, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto temp_mr = cudf::get_current_device_resource_ref(); @@ -422,7 +422,7 @@ assemble_build_column_info(shuffle_split_metadata const& h_global_metadata, // each partition linearly. to fix this, we'd have to change the kudo format in a way that would // increase it's size. I'm doing this as a kernel instead of through thrust so that I can // guarantee each partition is being marched by a seperate block to avoid thread divergence. - compute_offset_child_row_counts<<>>( + compute_offset_child_row_counts<<>>( offset_column_info, global_metadata, column_instance_info, @@ -652,7 +652,7 @@ rmm::device_uvector> transform_expand( SizeIterator first, SizeIterator last, GroupFunction op, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto temp_mr = cudf::get_current_device_resource_ref(); @@ -870,7 +870,7 @@ std::pair> assemble cudf::device_span partitions, cudf::device_span partition_offsets, size_t per_partition_metadata_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto temp_mr = cudf::get_current_device_resource_ref(); @@ -1348,7 +1348,7 @@ std::pair> assemble // Drain the async H2D uploads above before h_dst_buffers and the batched_memset span vector // (spans_to_zero) go out of scope: CUDA 13+ may read the host source only when the stream // executes the copy. - stream.synchronize(); + stream.sync(); // Return shuffle assemble result with slices and copy batches (column_views will be populated // later) @@ -1588,7 +1588,7 @@ __global__ void copy_offsets(cudf::device_span batches) void assemble_copy(cudf::device_span batches, cudf::device_span column_info, cudf::host_span h_column_info, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // TODO: it might make sense to launch these three copies on separate streams. It is likely that // the validity and offset copies will be much smaller than the data copies and could @@ -1636,12 +1636,12 @@ void assemble_copy(cudf::device_span batches, // copy validity constexpr int copy_validity_block_size = 128; copy_validity - <<>>(batches); + <<>>(batches); // copy offsets constexpr int copy_offsets_block_size = 128; copy_offsets - <<>>(batches); + <<>>(batches); // we have to sync because the build_table step will need the cpu-side valid_count in // h_column_info when constructing the columns. @@ -1650,7 +1650,7 @@ void assemble_copy(cudf::device_span batches, column_info.size() * sizeof(assemble_column_info), cudaMemcpyDefault, stream); - stream.synchronize(); + stream.sync(); } } // namespace @@ -1706,7 +1706,7 @@ calculate_empty_buffer_sizes(cudf::host_span col_i } // initialize buffers for empty columns -void initialize_empty_buffers(uint8_t* buffer_base, size_t total_size, rmm::cuda_stream_view stream) +void initialize_empty_buffers(uint8_t* buffer_base, size_t total_size, cuda::stream_ref stream) { // Initialize all buffers to 0, which is correct for: // - Validity: all bits 0 (but since num_rows=0, no bits matter) @@ -1752,7 +1752,7 @@ struct assemble_column_view_functor { cudf::host_span column_meta; cudf::host_span assemble_data; shuffle_assemble_result const& assemble_result; - rmm::cuda_stream_view stream; + cuda::stream_ref stream; rmm::device_async_resource_ref mr; template ())> @@ -1918,7 +1918,7 @@ struct assemble_column_view_functor { void build_table(cudf::host_span column_meta, cudf::host_span assemble_data, shuffle_assemble_result& assemble_result, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // create native cudf::column_view objects pointing to slices in the shared buffer @@ -1947,7 +1947,7 @@ void build_table(cudf::host_span column_meta, shuffle_assemble_result shuffle_assemble(shuffle_split_metadata const& metadata, cudf::device_span partitions, cudf::device_span _partition_offsets, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); diff --git a/src/main/cpp/src/shuffle_split.cu b/src/main/cpp/src/shuffle_split.cu index 88e4cd9c34..380ba38cb8 100644 --- a/src/main/cpp/src/shuffle_split.cu +++ b/src/main/cpp/src/shuffle_split.cu @@ -34,11 +34,11 @@ #include #include -#include #include #include #include +#include #include #include #include @@ -274,7 +274,7 @@ void setup_source_buf_info(InputIter begin, src_buf_info*& data_cur, int& offset_stack_pos, int& col_index, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, int parent_offset_index = -1, int offset_depth = 0); @@ -296,7 +296,7 @@ struct buf_info_functor { int& offset_stack_pos, int parent_offset_index, int offset_depth, - rmm::cuda_stream_view) + cuda::stream_ref) { if (include_nulls_for_column(col)) { add_null_buffer( @@ -354,7 +354,7 @@ void buf_info_functor::operator()(column_view const& col, int& offset_stack_pos, int parent_offset_index, int offset_depth, - rmm::cuda_stream_view) + cuda::stream_ref) { if (include_nulls_for_column(col)) { add_null_buffer( @@ -414,7 +414,7 @@ void buf_info_functor::operator()(column_view const& col, int& offset_stack_pos, int parent_offset_index, int offset_depth, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { lists_column_view lcv(col); @@ -479,7 +479,7 @@ void buf_info_functor::operator()(column_view const& col, int& offset_stack_pos, int parent_offset_index, int offset_depth, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (include_nulls_for_column(col)) { add_null_buffer( @@ -530,7 +530,7 @@ void setup_source_buf_info(InputIter begin, src_buf_info*& data_cur, int& offset_stack_pos, int& col_index, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, int parent_offset_index, int offset_depth) { @@ -708,7 +708,7 @@ void split_copy(src_buf_info const* src_bufs, uint8_t* dst_buf, size_t num_bufs, dst_buf_info const* d_dst_buf_info, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto input_iter = spark_rapids_jni::util::make_counting_transform_iterator( 0, cuda::proclaim_return_type([src_bufs, d_dst_buf_info] __device__(size_t i) { @@ -797,7 +797,7 @@ shuffle_split_metadata compute_metadata(cudf::table_view const& input, */ shuffle_split_output shuffle_split(cudf::table_view const& input, std::vector const& splits, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { SRJ_FUNC_RANGE(); @@ -890,7 +890,7 @@ shuffle_split_output shuffle_split(cudf::table_view const& input, // HtoD indices and source buf info to device CUDF_CUDA_TRY(cudaMemcpyAsync( - d_indices, h_indices, indices_size + src_buf_info_size, cudaMemcpyDefault, stream.value())); + d_indices, h_indices, indices_size + src_buf_info_size, cudaMemcpyDefault, stream.get())); // packed block of memory 2. partition buffer sizes, dst_buf_info structs and per-partition // has-validity buffer @@ -1136,7 +1136,7 @@ shuffle_split_output shuffle_split(cudf::table_view const& input, }); // allocate output buffer - stream.synchronize(); // for dst_buf_total_size from above + stream.sync(); // for dst_buf_total_size from above rmm::device_buffer dst_buf(dst_buf_total_size, stream, mr); // pack per-partition data. one thread per (flattened) column. @@ -1147,13 +1147,13 @@ shuffle_split_output shuffle_split(cudf::table_view const& input, pack_per_partition_metadata_kernel<<>>(static_cast(dst_buf.data()), - d_partition_offsets.data(), - num_partitions, - total_flattened_columns, - d_indices, - d_flattened_col_inst_has_validity, - d_partition_sizes); + stream.get()>>>(static_cast(dst_buf.data()), + d_partition_offsets.data(), + num_partitions, + total_flattened_columns, + d_indices, + d_flattened_col_inst_has_validity, + d_partition_sizes); // perform the copy. split_copy( @@ -1163,7 +1163,7 @@ shuffle_split_output shuffle_split(cudf::table_view const& input, // function only uses the cpu). auto metadata = compute_metadata(input, total_flattened_columns); - stream.synchronize(); + stream.sync(); return {shuffle_split_result{std::make_unique(std::move(dst_buf)), std::move(d_partition_offsets)}, std::move(metadata)}; diff --git a/src/main/cpp/src/shuffle_split.hpp b/src/main/cpp/src/shuffle_split.hpp index 946c46d8a1..7990a60b42 100644 --- a/src/main/cpp/src/shuffle_split.hpp +++ b/src/main/cpp/src/shuffle_split.hpp @@ -16,10 +16,10 @@ #include -#include #include #include +#include #include @@ -142,7 +142,7 @@ struct shuffle_split_output { */ shuffle_split_output shuffle_split(cudf::table_view const& input, std::vector const& splits, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -189,7 +189,7 @@ struct shuffle_assemble_result { shuffle_assemble_result shuffle_assemble(shuffle_split_metadata const& metadata, cudf::device_span partitions, cudf::device_span partition_offsets, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/substring_index.cu b/src/main/cpp/src/substring_index.cu index ee99890181..95c73f94ab 100644 --- a/src/main/cpp/src/substring_index.cu +++ b/src/main/cpp/src/substring_index.cu @@ -28,8 +28,7 @@ #include #include -#include - +#include #include #include #include @@ -55,7 +54,7 @@ void compute_substring_indices(column_device_view const& d_column, size_type delimiter_count, size_type* start_char_pos, size_type* end_char_pos, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref) { auto strings_count = d_column.size(); @@ -112,7 +111,7 @@ template std::unique_ptr substring_index(strings_column_view const& strings, DelimiterItrT const delimiter_itr, size_type count, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto strings_count = strings.size(); diff --git a/src/main/cpp/src/timezones.cu b/src/main/cpp/src/timezones.cu index 999f918845..b872635e22 100644 --- a/src/main/cpp/src/timezones.cu +++ b/src/main/cpp/src/timezones.cu @@ -34,13 +34,13 @@ #include #include -#include #include #include #include #include #include +#include #include using column = cudf::column; @@ -94,7 +94,7 @@ auto convert_timestamp_tz(column_view const& input, table_view const& transitions, size_type tz_index, bool to_utc, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // get the fixed transitions @@ -205,7 +205,7 @@ std::unique_ptr convert_to_utc_with_multiple_timezones( column_view const& tz_offset, table_view const& transitions, column_view const tz_indices, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input_seconds.type().id() == cudf::type_id::INT64, @@ -727,7 +727,7 @@ std::unique_ptr convert_timezones(cudf::column_view const& input, int64_t writer_2015_year_base_offset_us, spark_rapids_jni::orc_tz_side writer, spark_rapids_jni::orc_tz_side reader, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr, bool writer_reader_rules_differ) { @@ -794,7 +794,7 @@ std::unique_ptr convert_timezones(cudf::column_view const& input, auto const launch_config = cuda::make_config(cuda::grid_dims(num_blocks), cuda::block_dims(), cuda::dynamic_shared_memory(smem_bytes)); - cuda::launch(stream.value(), + cuda::launch(stream.get(), launch_config, convert_timezones_kernel, input.begin(), @@ -806,7 +806,7 @@ std::unique_ptr convert_timezones(cudf::column_view const& input, writer_args, reader_args, writer_reader_rules_differ); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); return results; } @@ -876,7 +876,7 @@ CUDF_KERNEL void __launch_bounds__(CONVERT_TZ_BLOCK_SIZE) template std::unique_ptr convert_orc_from_utc_typed(cudf::column_view const& input, spark_rapids_jni::orc_tz_side reader, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto results = cudf::make_fixed_width_column(input.type(), @@ -909,7 +909,7 @@ std::unique_ptr convert_orc_from_utc_typed(cudf::column_view const& inpu auto const launch_config = cuda::make_config(cuda::grid_dims(num_blocks), cuda::block_dims(), cuda::dynamic_shared_memory(smem_bytes)); - cuda::launch(stream.value(), + cuda::launch(stream.get(), launch_config, convert_orc_from_utc_kernel, input.begin(), @@ -918,7 +918,7 @@ std::unique_ptr convert_orc_from_utc_typed(cudf::column_view const& inpu input.size(), input.offset(), reader_args); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); return results; } @@ -932,7 +932,7 @@ std::unique_ptr convert_timestamp(column_view const& input, table_view const& transitions, size_type tz_index, bool to_utc, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const type = input.type().id(); @@ -960,7 +960,7 @@ std::unique_ptr convert_timestamp(column_view const& input, std::unique_ptr convert_timestamp_to_utc(column_view const& input, table_view const& transitions, size_type tz_index, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return convert_timestamp(input, transitions, tz_index, true, stream, mr); @@ -969,7 +969,7 @@ std::unique_ptr convert_timestamp_to_utc(column_view const& input, std::unique_ptr convert_utc_timestamp_to_timezone(column_view const& input, table_view const& transitions, size_type tz_index, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return convert_timestamp(input, transitions, tz_index, false, stream, mr); @@ -982,7 +982,7 @@ std::unique_ptr convert_timestamp_to_utc(column_view const& input_second column_view const& tz_offset, table_view const& transitions, column_view const tz_indices, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return convert_to_utc_with_multiple_timezones(input_seconds, @@ -1001,7 +1001,7 @@ std::unique_ptr convert_orc_writer_reader_timezones( int64_t writer_2015_year_base_offset_us, orc_tz_side writer, orc_tz_side reader, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr, bool writer_reader_rules_differ) { @@ -1011,7 +1011,7 @@ std::unique_ptr convert_orc_writer_reader_timezones( std::unique_ptr convert_orc_from_utc(cudf::column_view const& input, orc_tz_side reader, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { validate_timezone_table(reader.tz_info_table); @@ -1027,7 +1027,7 @@ std::unique_ptr convert_orc_writer_reader_timezones( int64_t writer_2015_year_base_offset_us, cudf::table_view const* reader_tz_info_table, int32_t reader_raw_offset, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // Java passes the exact ORC 2015 writer base offset, so this path does not infer it from diff --git a/src/main/cpp/src/timezones.hpp b/src/main/cpp/src/timezones.hpp index d91a5639b5..0fed427434 100644 --- a/src/main/cpp/src/timezones.hpp +++ b/src/main/cpp/src/timezones.hpp @@ -21,9 +21,10 @@ #include #include -#include #include +#include + #include #include #include @@ -49,7 +50,7 @@ std::unique_ptr convert_timestamp_to_utc( cudf::column_view const& input, cudf::table_view const& timezone_info, cudf::size_type const tz_index, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -71,7 +72,7 @@ std::unique_ptr convert_utc_timestamp_to_timezone( cudf::column_view const& input, cudf::table_view const& timezone_info, cudf::size_type const tz_index, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -102,7 +103,7 @@ std::unique_ptr convert_timestamp_to_utc( cudf::column_view const& tz_offset, cudf::table_view const& timezone_info, cudf::column_view const tz_indices, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -169,7 +170,7 @@ struct orc_tz_side { int64_t writer_2015_year_base_offset_us, orc_tz_side writer, orc_tz_side reader, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref(), bool writer_reader_rules_differ = true); @@ -186,7 +187,7 @@ struct orc_tz_side { [[nodiscard]] std::unique_ptr convert_orc_from_utc( cudf::column_view const& input, orc_tz_side reader, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -217,7 +218,7 @@ struct orc_tz_side { int64_t writer_2015_year_base_offset_us, cudf::table_view const* reader_tz_info_table, int32_t reader_raw_offset, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/utilities.cu b/src/main/cpp/src/utilities.cu index 93b234941c..a38edc0800 100644 --- a/src/main/cpp/src/utilities.cu +++ b/src/main/cpp/src/utilities.cu @@ -20,12 +20,12 @@ #include #include -#include #include #include #include #include +#include namespace spark_rapids_jni { @@ -38,7 +38,7 @@ bool is_basic_spark_numeric(cudf::data_type type) std::unique_ptr bitmask_bitwise_or( std::vector> const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input.size() > 0, "Empty input"); diff --git a/src/main/cpp/src/utilities.hpp b/src/main/cpp/src/utilities.hpp index 08376edefc..5d1e94863b 100644 --- a/src/main/cpp/src/utilities.hpp +++ b/src/main/cpp/src/utilities.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2025, NVIDIA CORPORATION. + * Copyright (c) 2023-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -22,10 +22,11 @@ #include #include -#include #include #include +#include + namespace spark_rapids_jni { template @@ -47,7 +48,7 @@ bool is_basic_spark_numeric(cudf::data_type type); */ std::unique_ptr bitmask_bitwise_or( std::vector> const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/uuid.cu b/src/main/cpp/src/uuid.cu index cc8d8960b0..fd339f158d 100644 --- a/src/main/cpp/src/uuid.cu +++ b/src/main/cpp/src/uuid.cu @@ -126,7 +126,7 @@ __launch_bounds__(block_size) CUDF_KERNEL void generate_uuids_kernel( std::unique_ptr generate_uuids(cudf::size_type row_count, long seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { @@ -170,7 +170,7 @@ std::unique_ptr generate_uuids(cudf::size_type row_count, rmm::device_uvector chars(num_chars, stream, mr); auto grid = cudf::detail::grid_1d(num_states, block_size); generate_uuids_kernel - <<>>( + <<>>( row_count, chars.data(), states.data(), num_states); return cudf::make_strings_column( @@ -186,7 +186,7 @@ std::unique_ptr generate_uuids(cudf::size_type row_count, std::unique_ptr random_uuids(int row_count, long seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { diff --git a/src/main/cpp/src/uuid.hpp b/src/main/cpp/src/uuid.hpp index 16ef46b872..809af0e147 100644 --- a/src/main/cpp/src/uuid.hpp +++ b/src/main/cpp/src/uuid.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2025, NVIDIA CORPORATION. + * Copyright (c) 2025-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -19,6 +19,8 @@ #include #include +#include + namespace spark_rapids_jni { /** @@ -39,7 +41,7 @@ namespace spark_rapids_jni { std::unique_ptr random_uuids( int row_count, long seed, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/src/zorder.cu b/src/main/cpp/src/zorder.cu index 7ac4d562e7..d649e60768 100644 --- a/src/main/cpp/src/zorder.cu +++ b/src/main/cpp/src/zorder.cu @@ -23,10 +23,10 @@ #include #include -#include #include #include +#include #include namespace { @@ -136,7 +136,7 @@ __device__ uint_backed_array hilbert_transposed_index( namespace spark_rapids_jni { std::unique_ptr interleave_bits(cudf::table_view const& tbl, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto num_columns = tbl.num_columns(); @@ -220,7 +220,7 @@ std::unique_ptr interleave_bits(cudf::table_view const& tbl, std::unique_ptr hilbert_index(int32_t const num_bits_per_entry, cudf::table_view const& tbl, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const num_rows = tbl.num_rows(); diff --git a/src/main/cpp/src/zorder.hpp b/src/main/cpp/src/zorder.hpp index 3e88f77ddd..1f10ea576e 100644 --- a/src/main/cpp/src/zorder.hpp +++ b/src/main/cpp/src/zorder.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2024, NVIDIA CORPORATION. + * Copyright (c) 2022-2026, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -19,22 +19,23 @@ #include #include -#include #include +#include + #include namespace spark_rapids_jni { std::unique_ptr interleave_bits( cudf::table_view const& tbl, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); std::unique_ptr hilbert_index( int32_t const num_bits, cudf::table_view const& tbl, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); } // namespace spark_rapids_jni diff --git a/src/main/cpp/tests/hyper_log_log_plus_plus.cu b/src/main/cpp/tests/hyper_log_log_plus_plus.cu index 31cbb4e228..ae9453a961 100644 --- a/src/main/cpp/tests/hyper_log_log_plus_plus.cu +++ b/src/main/cpp/tests/hyper_log_log_plus_plus.cu @@ -27,9 +27,10 @@ #include #include -#include #include +#include + #include #include #include @@ -86,7 +87,7 @@ std::vector get_column_ptrs_from_struct_scalars( std::unique_ptr make_struct_column_from_scalars( std::vector>& scalars, int num_longs_in_scalar, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // asserts @@ -123,7 +124,7 @@ std::unique_ptr make_struct_column_from_scalars( auto d_output = cudf::detail::make_device_uvector(host_results_pointers, stream, mr); // concatenate struct scalars into a struct column - concat_struct_scalars_to_struct_column_kernel<<<1, 1, 0, stream.value()>>>( + concat_struct_scalars_to_struct_column_kernel<<<1, 1, 0, stream.get()>>>( d_col_ptrs, scalars.size(), num_longs_in_scalar, d_output); // create struct column @@ -139,7 +140,7 @@ std::unique_ptr make_struct_column_from_scalars( */ std::unique_ptr make_struct_column_from_scalar(std::unique_ptr& scalar, int num_longs_in_scalar, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { std::vector> scalars; diff --git a/src/main/cpp/tests/protobuf_helpers.cu b/src/main/cpp/tests/protobuf_helpers.cu index 8097fc07aa..cb121403fd 100644 --- a/src/main/cpp/tests/protobuf_helpers.cu +++ b/src/main/cpp/tests/protobuf_helpers.cu @@ -26,6 +26,7 @@ #include +#include #include #include @@ -43,7 +44,7 @@ TEST_F(ProtobufHelpersTest, NullMaskFromPaddedValidUsesZeroLogicalRows) h_valid.data(), h_valid.size() * sizeof(h_valid[0]), cudaMemcpyDefault, - stream.value())); + stream.get())); auto [mask, null_count] = spark_rapids_jni::protobuf::detail::make_null_mask_from_valid( valid, 0, stream, cudf::get_current_device_resource_ref()); @@ -63,7 +64,7 @@ TEST_F(ProtobufHelpersTest, NullMaskFromPaddedValidIgnoresTail) h_valid.data(), h_valid.size() * sizeof(h_valid[0]), cudaMemcpyDefault, - stream.value())); + stream.get())); auto [mask, null_count] = spark_rapids_jni::protobuf::detail::make_null_mask_from_valid( valid, 2, stream, cudf::get_current_device_resource_ref()); @@ -73,8 +74,8 @@ TEST_F(ProtobufHelpersTest, NullMaskFromPaddedValidIgnoresTail) std::vector h_mask(mask.size() / sizeof(cudf::bitmask_type)); CUDF_CUDA_TRY( - cudaMemcpyAsync(h_mask.data(), mask.data(), mask.size(), cudaMemcpyDefault, stream.value())); - stream.synchronize(); + cudaMemcpyAsync(h_mask.data(), mask.data(), mask.size(), cudaMemcpyDefault, stream.get())); + stream.sync(); EXPECT_TRUE(cudf::bit_is_set(h_mask.data(), 0)); EXPECT_FALSE(cudf::bit_is_set(h_mask.data(), 1)); diff --git a/src/main/cpp/tests/shuffle_split.cu b/src/main/cpp/tests/shuffle_split.cu index 556bb171aa..e6d3b9dec5 100644 --- a/src/main/cpp/tests/shuffle_split.cu +++ b/src/main/cpp/tests/shuffle_split.cu @@ -49,7 +49,7 @@ spark_rapids_jni::shuffle_split_result reshape_partitions( cudf::device_span partitions, cudf::device_span partition_offsets, std::vector const& remaps, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(remaps.size() == partition_offsets.size() - 1, "Invaid remaps vector size"); From 62aa0a06ee1aced7bf05f2bef9ec099dd27a5bdb Mon Sep 17 00:00:00 2001 From: MithunR Date: Wed, 26 Aug 2026 16:52:33 -0700 Subject: [PATCH 2/6] Formatting. Signed-off-by: MithunR --- src/main/cpp/src/row_conversion.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/main/cpp/src/row_conversion.cu b/src/main/cpp/src/row_conversion.cu index 9d35bef5fa..f252809345 100644 --- a/src/main/cpp/src/row_conversion.cu +++ b/src/main/cpp/src/row_conversion.cu @@ -284,7 +284,7 @@ build_string_row_offsets(table_view const& tbl, */ struct string_row_offset_functor { string_row_offset_functor(device_span d_row_offsets) - : d_row_offsets(d_row_offsets){}; + : d_row_offsets(d_row_offsets) {}; __device__ inline size_type operator()(int row_number, int) const { @@ -300,7 +300,7 @@ struct string_row_offset_functor { */ struct fixed_width_row_offset_functor { fixed_width_row_offset_functor(size_type fixed_width_only_row_size) - : _fixed_width_only_row_size(fixed_width_only_row_size){}; + : _fixed_width_only_row_size(fixed_width_only_row_size) {}; __device__ inline size_type operator()(int row_number, int tile_row_start) const { From 8bf6b6da61ceddc97bf28ae52dafa3283c2d176d Mon Sep 17 00:00:00 2001 From: MithunR Date: Thu, 27 Aug 2026 16:02:59 -0700 Subject: [PATCH 3/6] Review changes: 1. Moved thirdparty/cudf to exclude NVIDIA/cudf/pull/23803. 2. Corrected kvikio version tag. Signed-off-by: MithunR --- thirdparty/cudf | 2 +- thirdparty/cudf-pins/versions.json | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 0137f3a6e2..40c1844e67 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 0137f3a6e2b9c4a1c4087030a9a5da6a814cde52 +Subproject commit 40c1844e6723bb45ef51314f40c1ed0d22a0a8a5 diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 03e5365be9..d78b5b1f6e 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -49,7 +49,7 @@ { "always_download" : true, "git_shallow" : false, - "git_tag" : "26bd4831585ce34294bce6bb3324d648cfd31747", + "git_tag" : "9af39e27161d680083b8c309b9040f8c944d68ad", "git_url" : "https://github.com/rapidsai/kvikio.git", "source_subdir" : "cpp", "version" : "26.10" From 4207de090bc35cb9b338b35b52bab9b72a64c0b0 Mon Sep 17 00:00:00 2001 From: MithunR Date: Mon, 31 Aug 2026 13:53:24 -0700 Subject: [PATCH 4/6] IWYU for . Signed-off-by: MithunR --- src/main/cpp/benchmarks/get_json_object.cu | 2 ++ src/main/cpp/src/aggregation64_utils.cu | 1 + src/main/cpp/src/case_when.cu | 1 + src/main/cpp/src/cast_string.cu | 1 + src/main/cpp/src/cast_string_to_float.cu | 1 + src/main/cpp/src/charset_decode.cu | 1 + src/main/cpp/src/datetime_rebase.cu | 1 + src/main/cpp/src/decimal_utils.cu | 1 + src/main/cpp/src/exception_with_row_index_utilities.cu | 1 + src/main/cpp/src/get_json_object.cu | 1 + src/main/cpp/src/histogram.cu | 1 + src/main/cpp/src/hyper_log_log_plus_plus.cu | 1 + src/main/cpp/src/hyper_log_log_plus_plus_host_udf.cu | 2 ++ src/main/cpp/src/iceberg/iceberg_bucket.cu | 1 + src/main/cpp/src/iceberg/iceberg_datetime_util.cu | 1 + src/main/cpp/src/iceberg/iceberg_truncate.cu | 1 + src/main/cpp/src/map.cu | 2 ++ src/main/cpp/src/map_zip_with_utils.cu | 1 + src/main/cpp/src/multiply.cu | 1 + src/main/cpp/src/number_converter.cu | 1 + src/main/cpp/src/protobuf/protobuf.cu | 1 + src/main/cpp/src/protobuf/protobuf_builders.cu | 1 + src/main/cpp/src/protobuf/protobuf_kernels.cu | 1 + src/main/cpp/src/uuid.cu | 1 + src/main/cpp/tests/shuffle_split.cu | 1 + 25 files changed, 28 insertions(+) diff --git a/src/main/cpp/benchmarks/get_json_object.cu b/src/main/cpp/benchmarks/get_json_object.cu index c78f6b321e..b523cc66ae 100644 --- a/src/main/cpp/benchmarks/get_json_object.cu +++ b/src/main/cpp/benchmarks/get_json_object.cu @@ -23,6 +23,8 @@ #include #include +#include + #include #include diff --git a/src/main/cpp/src/aggregation64_utils.cu b/src/main/cpp/src/aggregation64_utils.cu index 93613f404b..26bd62e90f 100644 --- a/src/main/cpp/src/aggregation64_utils.cu +++ b/src/main/cpp/src/aggregation64_utils.cu @@ -24,6 +24,7 @@ #include #include +#include #include #include diff --git a/src/main/cpp/src/case_when.cu b/src/main/cpp/src/case_when.cu index a15cfb531d..b99ddd4a11 100644 --- a/src/main/cpp/src/case_when.cu +++ b/src/main/cpp/src/case_when.cu @@ -22,6 +22,7 @@ #include #include +#include #include namespace spark_rapids_jni { diff --git a/src/main/cpp/src/cast_string.cu b/src/main/cpp/src/cast_string.cu index 1f35235ec4..aa1a3e14b1 100644 --- a/src/main/cpp/src/cast_string.cu +++ b/src/main/cpp/src/cast_string.cu @@ -33,6 +33,7 @@ #include #include #include +#include #include #include diff --git a/src/main/cpp/src/cast_string_to_float.cu b/src/main/cpp/src/cast_string_to_float.cu index 42461bbd3a..9a671efd61 100644 --- a/src/main/cpp/src/cast_string_to_float.cu +++ b/src/main/cpp/src/cast_string_to_float.cu @@ -34,6 +34,7 @@ #include #include #include +#include using namespace cudf; diff --git a/src/main/cpp/src/charset_decode.cu b/src/main/cpp/src/charset_decode.cu index e42d4bfb9b..1e5a9c7af7 100644 --- a/src/main/cpp/src/charset_decode.cu +++ b/src/main/cpp/src/charset_decode.cu @@ -26,6 +26,7 @@ #include #include +#include #include diff --git a/src/main/cpp/src/datetime_rebase.cu b/src/main/cpp/src/datetime_rebase.cu index 3955f5c078..c8e0799186 100644 --- a/src/main/cpp/src/datetime_rebase.cu +++ b/src/main/cpp/src/datetime_rebase.cu @@ -26,6 +26,7 @@ #include #include +#include #include #include diff --git a/src/main/cpp/src/decimal_utils.cu b/src/main/cpp/src/decimal_utils.cu index c3a444078c..08e23e21a8 100644 --- a/src/main/cpp/src/decimal_utils.cu +++ b/src/main/cpp/src/decimal_utils.cu @@ -33,6 +33,7 @@ #include #include +#include #include #include diff --git a/src/main/cpp/src/exception_with_row_index_utilities.cu b/src/main/cpp/src/exception_with_row_index_utilities.cu index eaac777bc7..554cf869d6 100644 --- a/src/main/cpp/src/exception_with_row_index_utilities.cu +++ b/src/main/cpp/src/exception_with_row_index_utilities.cu @@ -22,6 +22,7 @@ #include +#include #include #include diff --git a/src/main/cpp/src/get_json_object.cu b/src/main/cpp/src/get_json_object.cu index 4907b9f4a6..699ad4ccb1 100644 --- a/src/main/cpp/src/get_json_object.cu +++ b/src/main/cpp/src/get_json_object.cu @@ -41,6 +41,7 @@ #include #include #include +#include #include #include diff --git a/src/main/cpp/src/histogram.cu b/src/main/cpp/src/histogram.cu index fbc3d064ed..4591099bdc 100644 --- a/src/main/cpp/src/histogram.cu +++ b/src/main/cpp/src/histogram.cu @@ -34,6 +34,7 @@ #include #include +#include #include #include #include diff --git a/src/main/cpp/src/hyper_log_log_plus_plus.cu b/src/main/cpp/src/hyper_log_log_plus_plus.cu index 992555d598..964d66f7fd 100644 --- a/src/main/cpp/src/hyper_log_log_plus_plus.cu +++ b/src/main/cpp/src/hyper_log_log_plus_plus.cu @@ -39,6 +39,7 @@ #include #include // TODO #include once available #include +#include #include #include diff --git a/src/main/cpp/src/hyper_log_log_plus_plus_host_udf.cu b/src/main/cpp/src/hyper_log_log_plus_plus_host_udf.cu index f4ce0e4392..12402e33a6 100644 --- a/src/main/cpp/src/hyper_log_log_plus_plus_host_udf.cu +++ b/src/main/cpp/src/hyper_log_log_plus_plus_host_udf.cu @@ -22,6 +22,8 @@ #include #include +#include + namespace spark_rapids_jni { namespace { diff --git a/src/main/cpp/src/iceberg/iceberg_bucket.cu b/src/main/cpp/src/iceberg/iceberg_bucket.cu index f486f8e785..411397d0ae 100644 --- a/src/main/cpp/src/iceberg/iceberg_bucket.cu +++ b/src/main/cpp/src/iceberg/iceberg_bucket.cu @@ -35,6 +35,7 @@ #include #include #include +#include #include #include diff --git a/src/main/cpp/src/iceberg/iceberg_datetime_util.cu b/src/main/cpp/src/iceberg/iceberg_datetime_util.cu index ff0a213216..f16defcad3 100644 --- a/src/main/cpp/src/iceberg/iceberg_datetime_util.cu +++ b/src/main/cpp/src/iceberg/iceberg_datetime_util.cu @@ -26,6 +26,7 @@ #include +#include #include namespace spark_rapids_jni { diff --git a/src/main/cpp/src/iceberg/iceberg_truncate.cu b/src/main/cpp/src/iceberg/iceberg_truncate.cu index eff3e9bc10..140423dfba 100644 --- a/src/main/cpp/src/iceberg/iceberg_truncate.cu +++ b/src/main/cpp/src/iceberg/iceberg_truncate.cu @@ -30,6 +30,7 @@ #include #include +#include #include #include diff --git a/src/main/cpp/src/map.cu b/src/main/cpp/src/map.cu index 2efe5c595a..6e8f514e9d 100644 --- a/src/main/cpp/src/map.cu +++ b/src/main/cpp/src/map.cu @@ -24,6 +24,8 @@ #include #include +#include + namespace spark_rapids_jni { std::unique_ptr sort_map_column(cudf::column_view const& input, diff --git a/src/main/cpp/src/map_zip_with_utils.cu b/src/main/cpp/src/map_zip_with_utils.cu index 226407e272..d1b3697e3d 100644 --- a/src/main/cpp/src/map_zip_with_utils.cu +++ b/src/main/cpp/src/map_zip_with_utils.cu @@ -29,6 +29,7 @@ #include #include +#include #include using namespace cudf; diff --git a/src/main/cpp/src/multiply.cu b/src/main/cpp/src/multiply.cu index 08e9154c95..810b1629dd 100644 --- a/src/main/cpp/src/multiply.cu +++ b/src/main/cpp/src/multiply.cu @@ -29,6 +29,7 @@ #include #include +#include #include #include diff --git a/src/main/cpp/src/number_converter.cu b/src/main/cpp/src/number_converter.cu index 5833aefe07..4c0f238216 100644 --- a/src/main/cpp/src/number_converter.cu +++ b/src/main/cpp/src/number_converter.cu @@ -27,6 +27,7 @@ #include #include #include +#include #include #include diff --git a/src/main/cpp/src/protobuf/protobuf.cu b/src/main/cpp/src/protobuf/protobuf.cu index f6b046e151..47e6865d8c 100644 --- a/src/main/cpp/src/protobuf/protobuf.cu +++ b/src/main/cpp/src/protobuf/protobuf.cu @@ -23,6 +23,7 @@ #include #include +#include #include #include diff --git a/src/main/cpp/src/protobuf/protobuf_builders.cu b/src/main/cpp/src/protobuf/protobuf_builders.cu index 1b8109d709..45e407201d 100644 --- a/src/main/cpp/src/protobuf/protobuf_builders.cu +++ b/src/main/cpp/src/protobuf/protobuf_builders.cu @@ -19,6 +19,7 @@ #include #include +#include #include #include diff --git a/src/main/cpp/src/protobuf/protobuf_kernels.cu b/src/main/cpp/src/protobuf/protobuf_kernels.cu index f6acca9fb8..7d03326de6 100644 --- a/src/main/cpp/src/protobuf/protobuf_kernels.cu +++ b/src/main/cpp/src/protobuf/protobuf_kernels.cu @@ -23,6 +23,7 @@ #include #include +#include #include #include #include diff --git a/src/main/cpp/src/uuid.cu b/src/main/cpp/src/uuid.cu index fd339f158d..0f98cae8ec 100644 --- a/src/main/cpp/src/uuid.cu +++ b/src/main/cpp/src/uuid.cu @@ -27,6 +27,7 @@ #include #include +#include #include #include diff --git a/src/main/cpp/tests/shuffle_split.cu b/src/main/cpp/tests/shuffle_split.cu index 1773acdc4d..d5e9703a02 100644 --- a/src/main/cpp/tests/shuffle_split.cu +++ b/src/main/cpp/tests/shuffle_split.cu @@ -29,6 +29,7 @@ #include #include +#include struct ShuffleSplitTests : public cudf::test::BaseFixture {}; From a7a33aa87b6e8e10edfcf4f455df91084f286146 Mon Sep 17 00:00:00 2001 From: MithunR Date: Mon, 31 Aug 2026 14:17:17 -0700 Subject: [PATCH 5/6] Review: Changed default arg in regex_rewrite_utils. Signed-off-by: MithunR --- src/main/cpp/src/regex_rewrite_utils.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/main/cpp/src/regex_rewrite_utils.hpp b/src/main/cpp/src/regex_rewrite_utils.hpp index 8b83ec7d10..26668b08e4 100644 --- a/src/main/cpp/src/regex_rewrite_utils.hpp +++ b/src/main/cpp/src/regex_rewrite_utils.hpp @@ -42,6 +42,6 @@ std::unique_ptr literal_range_pattern( int const len, int const start, int const end, - cuda::stream_ref stream = rmm::cuda_stream_default, + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()); } // namespace spark_rapids_jni From afa5b99bf761550cceb136770ed52085875d2c57 Mon Sep 17 00:00:00 2001 From: MithunR Date: Mon, 31 Aug 2026 14:26:45 -0700 Subject: [PATCH 6/6] Latest cudf commit. Signed-off-by: MithunR --- thirdparty/cudf | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thirdparty/cudf b/thirdparty/cudf index 75ec3b97e5..7385730e6d 160000 --- a/thirdparty/cudf +++ b/thirdparty/cudf @@ -1 +1 @@ -Subproject commit 75ec3b97e5f544b4bf9427484f8a52134cc71ecb +Subproject commit 7385730e6d2eb9c0c9d572e9138f6335e800cbba