diff --git a/src/main/cpp/benchmarks/bloom_filter.cu b/src/main/cpp/benchmarks/bloom_filter.cu index 9f48c838d9..8a04293db8 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 @@ -39,13 +41,13 @@ void bloom_filter_put_impl(nvbench::state& state, int version) auto const src = create_random_table({{cudf::type_id::INT64}}, row_count{num_rows}, builder); 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())); + cuda::stream_ref const stream = cudf::get_default_stream(); + 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..4f846f666f 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 @@ -27,10 +29,10 @@ static void long_to_binary_string(nvbench::state& state) { auto const num_rows = static_cast(state.get_int64("num_rows")); - 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())); + auto const input_table = create_random_table({cudf::type_id::INT64}, row_count{num_rows}); + auto const long_col = input_table->get_column(0); + cuda::stream_ref const stream = cudf::get_default_stream(); + 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..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 @@ -44,7 +46,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 +64,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..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 @@ -75,7 +76,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 +113,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..4e801391b0 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) { @@ -162,14 +162,14 @@ void pack_bloom_filter_header(cudf::device_span buf, byte_swap_int32(header.num_hashes), byte_swap_int32(header.num_longs)}; CUDF_CUDA_TRY(cudaMemcpyAsync( - buf.data(), &raw, bloom_filter_header_v1_size_bytes, cudaMemcpyDefault, stream)); + buf.data(), &raw, bloom_filter_header_v1_size_bytes, cudaMemcpyDefault, stream.get())); } else { bloom_filter_header_v2 raw = {byte_swap_int32(header.version), byte_swap_int32(header.num_hashes), byte_swap_int32(seed), byte_swap_int32(header.num_longs)}; CUDF_CUDA_TRY(cudaMemcpyAsync( - buf.data(), &raw, bloom_filter_header_v2_size_bytes, cudaMemcpyDefault, stream)); + buf.data(), &raw, bloom_filter_header_v2_size_bytes, cudaMemcpyDefault, stream.get())); } } @@ -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"); @@ -200,8 +200,8 @@ unpack_bloom_filter(cudf::device_span bloom_filter, rmm::cuda_str // TODO (future): Consider using pinned host memory for cudaMemcpyAsync. // 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(); + cudaMemcpyAsync(raw_ints, bloom_filter.data(), read_size, cudaMemcpyDefault, stream.get())); + 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(); @@ -320,8 +320,8 @@ std::unique_ptr bloom_filter_create(int version, stream, (version == bloom_filter_version_1 ? 0 : seed)); - CUDF_CUDA_TRY( - cudaMemsetAsync(static_cast(buf.data()) + hdr_size, 0, bloom_filter_size, stream)); + CUDF_CUDA_TRY(cudaMemsetAsync( + static_cast(buf.data()) + hdr_size, 0, bloom_filter_size, stream.get())); return std::make_unique( cudf::column( @@ -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..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 { @@ -67,7 +68,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 +96,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..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 @@ -610,7 +611,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 +628,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 +638,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 +660,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 +675,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 +705,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 +732,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 +744,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 +774,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 +799,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 +824,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 a6ab8fb1b0..ba31e29d0d 100644 --- a/src/main/cpp/src/cast_string.hpp +++ b/src/main/cpp/src/cast_string.hpp @@ -25,6 +25,8 @@ #include +#include + #include #include @@ -79,7 +81,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()); /** @@ -101,7 +103,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()); /** @@ -119,34 +121,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()); /** @@ -177,7 +179,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()); /** @@ -199,7 +201,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()); /** @@ -245,7 +247,7 @@ std::unique_ptr parse_timestamp_strings_with_format( std::string const& format, bool legacy, bool exception_policy, - 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()); /** @@ -261,7 +263,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..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; @@ -886,7 +887,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 +942,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 +952,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 b37f54aa18..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 @@ -202,7 +203,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, @@ -257,7 +258,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..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 @@ -53,7 +54,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 +125,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 +225,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 +288,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 +342,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 +358,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 0850d201b7..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 @@ -969,7 +970,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 +1005,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 +1039,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 +1072,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 +1105,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 +1138,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 +1366,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 +1389,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..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 @@ -91,7 +92,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 +116,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 +139,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 +163,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 +181,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..2e687d51bd 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] = @@ -192,7 +192,7 @@ std::tuple> unify_json_s concatenated_buff->data(), concatenated_buff->size(), cudaMemcpyDefault, - stream)); + stream.get())); cudf::detail::cuda_memcpy_async( cudf::device_span(static_cast(unified_buff.data()) + concatenated_buff->size(), 1u), @@ -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..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 @@ -870,7 +871,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 +886,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 +907,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 +967,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 +1019,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 +1174,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 +1248,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 +1261,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 307b7b5c70..cd978ffde7 100644 --- a/src/main/cpp/src/hash/hash.hpp +++ b/src/main/cpp/src/hash/hash.hpp @@ -19,10 +19,9 @@ #include #include -#include #include -#include +#include namespace spark_rapids_jni { @@ -42,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()); /** @@ -58,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()); /** @@ -72,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()); /** 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 5eb0512110..ec7036973d 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 6b095bbdad..f82bb88869 100644 --- a/src/main/cpp/src/hash/sha.cpp +++ b/src/main/cpp/src/hash/sha.cpp @@ -21,7 +21,7 @@ #include -#include +#include #include 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..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 @@ -164,7 +165,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 +256,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 +276,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 +417,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.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.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/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 d2bcfe0849..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 @@ -383,7 +384,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 +394,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 +468,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..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 { @@ -136,7 +137,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 +166,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 +195,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 +226,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 +252,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 +260,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 +268,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 +276,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..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 @@ -131,7 +132,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 +142,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 +171,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 +199,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 +244,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 +253,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 +262,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..6e8f514e9d 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. @@ -24,11 +24,13 @@ #include #include +#include + 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 903b2a921c..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; @@ -50,7 +51,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 +113,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 +310,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..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 @@ -181,7 +182,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 +240,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 +308,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 +321,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 +345,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 +369,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..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 @@ -368,7 +369,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 +455,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 +480,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 +628,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 2f29a2b73b..5b45d16d42 100644 --- a/src/main/cpp/src/parse_timestamp_with_format.cu +++ b/src/main/cpp/src/parse_timestamp_with_format.cu @@ -31,12 +31,12 @@ #include #include -#include #include #include #include #include +#include #include #include @@ -361,7 +361,7 @@ std::unique_ptr parse_timestamp_strings_with_format( std::string const& format, bool legacy, bool exception_policy, - 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.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 23d4650a49..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 @@ -41,7 +42,7 @@ namespace detail { 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]; @@ -293,7 +294,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}; @@ -374,7 +375,7 @@ std::unique_ptr decode_protobuf_to_struct(cudf::column_view const& track_permissive_null_rows ? num_rows : 0, stream, scratch_mr); 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}; auto const recursive_context = recursive_decode_context{schema_context, decode_ctx}; @@ -519,8 +520,8 @@ 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); - CUDF_CHECK_CUDA(stream.value()); + <<>>(batch_input); + CUDF_CHECK_CUDA(stream.get()); } for (int j = 0; j < nf; j++) { @@ -721,11 +722,11 @@ std::unique_ptr decode_protobuf_to_struct(cudf::column_view const& { using enum protobuf_error; - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); 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)); } @@ -766,7 +767,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 8f7ece66c3..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 @@ -35,7 +36,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 +59,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 +107,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 +123,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 +140,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); } @@ -181,7 +182,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()) { @@ -193,7 +194,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( @@ -211,7 +212,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); @@ -223,7 +224,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}, @@ -231,8 +232,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{}); } @@ -257,7 +258,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) { auto const scratch_mr = cudf::get_current_device_resource_ref(); auto d_valid_enums = cudf::detail::make_device_uvector_async(valid_enums, stream, scratch_mr); @@ -300,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) { auto const scratch_mr = cudf::get_current_device_resource_ref(); @@ -325,7 +326,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 +343,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); @@ -383,7 +384,7 @@ std::unique_ptr build_repeated_string_column(cudf::column_view con protobuf_input_view input, repeated_field_work work, bool is_bytes, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { validate_nonempty_repeated_field_work(work, input.num_rows); @@ -397,8 +398,8 @@ std::unique_ptr build_repeated_string_column(cudf::column_view con repeated_location_provider loc_provider{ input.row_offsets, input.base_offset, work.occurrences->data()}; extract_lengths_kernel - <<>>(loc_provider, total_count, str_lengths.data()); - CUDF_CHECK_CUDA(stream.value()); + <<>>(loc_provider, total_count, str_lengths.data()); + CUDF_CHECK_CUDA(stream.get()); auto [str_offsets_col, total_chars] = cudf::strings::detail::make_offsets_child_column( str_lengths.begin(), str_lengths.end(), stream, mr); @@ -434,7 +435,7 @@ std::unique_ptr build_repeated_string_column(cudf::column_view con size_t temp_storage_bytes = 0; CUDF_CUDA_TRY(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); CUDF_CUDA_TRY(cub::DeviceMemcpy::Batched(temp_storage.data(), temp_storage_bytes, @@ -442,7 +443,7 @@ std::unique_ptr build_repeated_string_column(cudf::column_view con dst_iter, size_iter, total_count, - stream.value())); + stream.get())); } std::unique_ptr child_col; @@ -481,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; @@ -621,7 +622,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 c38ce01ec7..37dafc4e62 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,14 +496,14 @@ 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(cudf::column_view const& binary_input, protobuf_input_view input, repeated_field_work work, bool is_bytes, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr build_nested_struct_column( @@ -515,14 +512,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 130de36520..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 @@ -732,53 +733,51 @@ 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); - CUDF_CHECK_CUDA(stream.value()); + set_error_if_unset_kernel<<<1, 1, 0, stream.get()>>>(error_flag, error); + CUDF_CHECK_CUDA(stream.get()); } 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); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); } 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); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); } 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); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); } void launch_extract_strided_locations(field_location const* nested_locations, @@ -786,13 +785,13 @@ 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); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); } void launch_scan_nested_message_fields(protobuf_input_view input, @@ -800,85 +799,85 @@ 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); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); } 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); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); } void launch_compute_grandchild_parent_locations(nested_location_provider loc_provider, 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<<>>( loc_provider, gc_parent_locs, num_rows, error_flag); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); } 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); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); } 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); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); } 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); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); } 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; } @@ -897,19 +896,19 @@ 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()), !decode_ctx.row_force_null->is_empty() ? decode_ctx.row_force_null->data() : nullptr, decode_ctx.error->data()); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); } 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; @@ -955,7 +954,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; @@ -974,7 +973,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 073f2c3073..32e1ac1a42 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 @@ -390,38 +391,38 @@ 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(nested_location_provider loc_provider, 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 @@ -432,7 +433,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"); @@ -456,14 +457,14 @@ 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); - CUDF_CHECK_CUDA(stream.value()); + extract_scalar_kernel + <<>>(message_data, loc_provider, num_rows, output, options); + CUDF_CHECK_CUDA(stream.get()); }); } @@ -488,7 +489,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); } @@ -517,7 +518,7 @@ inline std::unique_ptr extract_and_build_string_or_bytes_column( ValidityFn validity_fn, bool has_default, cudf::detail::host_vector const& default_bytes, - 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; @@ -530,9 +531,9 @@ inline std::unique_ptr extract_and_build_string_or_bytes_column( rmm::device_uvector lengths(num_rows, stream, scratch_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); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); auto [offsets_col, total_size] = cudf::strings::detail::make_offsets_child_column(lengths.begin(), lengths.end(), stream, mr); @@ -573,7 +574,7 @@ inline std::unique_ptr extract_and_build_string_or_bytes_column( size_t temp_storage_bytes = 0; CUDF_CUDA_TRY(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, scratch_mr); CUDF_CUDA_TRY(cub::DeviceMemcpy::Batched(temp_storage.data(), temp_storage_bytes, @@ -581,7 +582,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) { @@ -616,7 +617,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); @@ -676,7 +677,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; @@ -755,7 +756,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); @@ -788,22 +789,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..26668b08e4 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 = 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/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..c2a5f44e27 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 @@ -1620,8 +1620,13 @@ void assemble_copy(cudf::device_span batches, })); size_t temp_storage_bytes{0}; - cub::DeviceMemcpy::Batched( - nullptr, temp_storage_bytes, input_iter, output_iter, size_iter, batches.size(), stream); + cub::DeviceMemcpy::Batched(nullptr, + temp_storage_bytes, + input_iter, + output_iter, + size_iter, + batches.size(), + stream.get()); rmm::device_buffer temp_storage( temp_storage_bytes, stream, cudf::get_current_device_resource_ref()); cub::DeviceMemcpy::Batched(temp_storage.data(), @@ -1630,18 +1635,18 @@ void assemble_copy(cudf::device_span batches, output_iter, size_iter, batches.size(), - stream); + stream.get()); } // 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. @@ -1649,8 +1654,8 @@ void assemble_copy(cudf::device_span batches, column_info.data(), column_info.size() * sizeof(assemble_column_info), cudaMemcpyDefault, - stream); - stream.synchronize(); + stream.get()); + stream.sync(); } } // namespace @@ -1706,13 +1711,13 @@ 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) // - Offsets: 0 (correct for empty offsets) // - Data: 0 (empty data) - cudaMemsetAsync(buffer_base, 0, total_size, stream); + cudaMemsetAsync(buffer_base, 0, total_size, stream.get()); } // create buffer slices for empty columns @@ -1752,7 +1757,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 +1923,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 +1952,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..8a0338dd45 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) { @@ -729,11 +729,16 @@ void split_copy(src_buf_info const* src_bufs, size_t temp_storage_bytes = 0; // Initialized on the next line, not that the compiler would know. cub::DeviceMemcpy::Batched( - nullptr, temp_storage_bytes, input_iter, output_iter, size_iter, num_bufs, stream); + nullptr, temp_storage_bytes, input_iter, output_iter, size_iter, num_bufs, stream.get()); rmm::device_buffer temp_storage( temp_storage_bytes, stream, cudf::get_current_device_resource_ref()); - cub::DeviceMemcpy::Batched( - temp_storage.data(), temp_storage_bytes, input_iter, output_iter, size_iter, num_bufs, stream); + cub::DeviceMemcpy::Batched(temp_storage.data(), + temp_storage_bytes, + input_iter, + output_iter, + size_iter, + num_bufs, + stream.get()); } /** @@ -797,7 +802,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 +895,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 @@ -1082,7 +1087,7 @@ shuffle_split_output shuffle_split(cudf::table_view const& input, d_partition_offsets.begin() + num_partitions, sizeof(size_t), cudaMemcpyDefault, - stream); + stream.get()); // generate destination offsets for each of the source copies, by partition, by section. auto buf_sizes = spark_rapids_jni::util::make_counting_transform_iterator( @@ -1136,7 +1141,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 +1152,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 +1168,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..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 @@ -126,7 +127,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 +171,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 +187,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 1bd288d413..7fa6a53371 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 @@ -35,7 +36,7 @@ class ProtobufHelpersTest : public cudf::test::BaseFixture {}; TEST_F(ProtobufHelpersTest, NullMaskFromPaddedValidUsesZeroLogicalRows) { - auto stream = cudf::get_default_stream(); + cuda::stream_ref stream = cudf::get_default_stream(); std::array h_valid{false}; rmm::device_uvector valid(h_valid.size(), stream); @@ -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()); @@ -55,7 +56,7 @@ TEST_F(ProtobufHelpersTest, NullMaskFromPaddedValidUsesZeroLogicalRows) TEST_F(ProtobufHelpersTest, NullMaskFromPaddedValidIgnoresTail) { - auto stream = cudf::get_default_stream(); + cuda::stream_ref stream = cudf::get_default_stream(); std::array h_valid{true, false, false}; rmm::device_uvector valid(h_valid.size(), stream); @@ -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)); @@ -82,7 +83,7 @@ TEST_F(ProtobufHelpersTest, NullMaskFromPaddedValidIgnoresTail) TEST_F(ProtobufHelpersTest, NullMaskFromAllValidRowsIsEmpty) { - auto stream = cudf::get_default_stream(); + cuda::stream_ref stream = cudf::get_default_stream(); std::array h_valid{true, true}; rmm::device_uvector valid(h_valid.size(), stream); @@ -90,7 +91,7 @@ TEST_F(ProtobufHelpersTest, NullMaskFromAllValidRowsIsEmpty) 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, h_valid.size(), stream, cudf::get_current_device_resource_ref()); diff --git a/src/main/cpp/tests/shuffle_split.cu b/src/main/cpp/tests/shuffle_split.cu index 556bb171aa..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 {}; @@ -49,7 +50,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"); @@ -100,7 +101,7 @@ spark_rapids_jni::shuffle_split_result reshape_partitions( size_t temp_storage_bytes = 0; cub::DeviceMemcpy::Batched( - nullptr, temp_storage_bytes, input_iter, output_iter, size_iter, num_partitions, stream); + nullptr, temp_storage_bytes, input_iter, output_iter, size_iter, num_partitions, stream.get()); rmm::device_buffer temp_storage( temp_storage_bytes, stream, cudf::get_current_device_resource_ref()); cub::DeviceMemcpy::Batched(temp_storage.data(), @@ -109,7 +110,7 @@ spark_rapids_jni::shuffle_split_result reshape_partitions( output_iter, size_iter, num_partitions, - stream); + stream.get()); return {std::make_unique(std::move(remapped_partitions)), std::move(remapped_offsets)}; diff --git a/thirdparty/cudf-pins/versions.json b/thirdparty/cudf-pins/versions.json index 36060364cd..c71f692072 100644 --- a/thirdparty/cudf-pins/versions.json +++ b/thirdparty/cudf-pins/versions.json @@ -120,4 +120,4 @@ "version" : "1.5.7" } } -} \ No newline at end of file +}