diff --git a/cpp/include/cudf/ast/detail/expression_evaluator.cuh b/cpp/include/cudf/ast/detail/expression_evaluator.cuh index 1a95bf16382f..5e6539d74e18 100644 --- a/cpp/include/cudf/ast/detail/expression_evaluator.cuh +++ b/cpp/include/cudf/ast/detail/expression_evaluator.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -16,10 +16,9 @@ #include #include -#include - #include #include +#include namespace cudf::ast::detail { diff --git a/cpp/include/cudf/ast/detail/expression_parser.hpp b/cpp/include/cudf/ast/detail/expression_parser.hpp index ef646c1571b4..37e922ec57d2 100644 --- a/cpp/include/cudf/ast/detail/expression_parser.hpp +++ b/cpp/include/cudf/ast/detail/expression_parser.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -106,7 +106,7 @@ class expression_parser { cudf::table_view const& left, std::optional> right, bool has_nulls, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -121,7 +121,7 @@ class expression_parser { expression_parser(expression const& expr, cudf::table_view const& table, bool has_nulls, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -240,7 +240,7 @@ class expression_parser { alignment = std::max(alignment, static_cast(alignof(T))); } - void move_to_device(rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr); + void move_to_device(cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @brief Helper function for recursive traversal of expressions. diff --git a/cpp/include/cudf/ast/expressions.hpp b/cpp/include/cudf/ast/expressions.hpp index 4329f6904300..f259082c950e 100644 --- a/cpp/include/cudf/ast/expressions.hpp +++ b/cpp/include/cudf/ast/expressions.hpp @@ -13,6 +13,8 @@ #include #include +#include + #include #include #include @@ -101,7 +103,7 @@ struct [[nodiscard]] expression { * @param stream CUDA stream used for device memory operations and kernel launches * @return `true` if the expression may evaluate to null, otherwise false */ - [[nodiscard]] bool may_evaluate_null(table_view const& left, rmm::cuda_stream_view stream) const + [[nodiscard]] bool may_evaluate_null(table_view const& left, cuda::stream_ref stream) const { return may_evaluate_null(left, left, stream); } @@ -116,7 +118,7 @@ struct [[nodiscard]] expression { */ [[nodiscard]] virtual bool may_evaluate_null(table_view const& left, table_view const& right, - rmm::cuda_stream_view stream) const = 0; + cuda::stream_ref stream) const = 0; virtual ~expression() {} }; @@ -384,7 +386,7 @@ class literal : public expression { [[nodiscard]] bool may_evaluate_null(table_view const& left, table_view const& right, - rmm::cuda_stream_view stream) const override + cuda::stream_ref stream) const override { return !is_valid(stream); } @@ -395,7 +397,7 @@ class literal : public expression { * @param stream CUDA stream used for device memory operations and kernel launches * @return true if the underlying scalar is valid */ - [[nodiscard]] bool is_valid(rmm::cuda_stream_view stream) const + [[nodiscard]] bool is_valid(cuda::stream_ref stream) const { if (auto* s = std::get_if(&scalar)) { return s->scalar.get().is_valid(stream); @@ -477,7 +479,7 @@ class column_reference : public expression { [[nodiscard]] bool may_evaluate_null(table_view const& left, table_view const& right, - rmm::cuda_stream_view stream) const override + cuda::stream_ref stream) const override { return (table_source == table_reference::LEFT ? left : right).column(column_index).has_nulls(); } @@ -552,7 +554,7 @@ class operation : public expression { [[nodiscard]] bool may_evaluate_null(table_view const& left, table_view const& right, - rmm::cuda_stream_view stream) const override; + cuda::stream_ref stream) const override; /** * @copydoc expression::accept @@ -592,7 +594,7 @@ class predicate : public expression { [[nodiscard]] bool may_evaluate_null(table_view const& left, table_view const& right, - rmm::cuda_stream_view stream) const override; + cuda::stream_ref stream) const override; /** * @copydoc expression::accept @@ -645,7 +647,7 @@ class column_name_reference : public expression { [[nodiscard]] bool may_evaluate_null(table_view const& left, table_view const& right, - rmm::cuda_stream_view stream) const override + cuda::stream_ref stream) const override { return true; } diff --git a/cpp/include/cudf/datetime.hpp b/cpp/include/cudf/datetime.hpp index 80ef6fb4639d..7a1083469c0b 100644 --- a/cpp/include/cudf/datetime.hpp +++ b/cpp/include/cudf/datetime.hpp @@ -10,7 +10,7 @@ #include #include -#include +#include #include @@ -57,7 +57,7 @@ enum class datetime_component : uint8_t { std::unique_ptr extract_datetime_component( cudf::column_view const& column, datetime_component component, - 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()); /** @} */ // end of group @@ -79,7 +79,7 @@ std::unique_ptr extract_datetime_component( */ std::unique_ptr last_day_of_month( cudf::column_view const& column, - 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()); /** @@ -95,7 +95,7 @@ std::unique_ptr last_day_of_month( */ std::unique_ptr day_of_year( cudf::column_view const& column, - 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()); /** @@ -132,7 +132,7 @@ std::unique_ptr day_of_year( std::unique_ptr add_calendrical_months( cudf::column_view const& timestamps, cudf::column_view const& months, - 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 +169,7 @@ std::unique_ptr add_calendrical_months( std::unique_ptr add_calendrical_months( cudf::column_view const& timestamps, cudf::scalar const& months, - 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()); /** @@ -188,7 +188,7 @@ std::unique_ptr add_calendrical_months( */ std::unique_ptr is_leap_year( cudf::column_view const& column, - 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()); /** @@ -207,7 +207,7 @@ std::unique_ptr is_leap_year( */ std::unique_ptr days_in_month( cudf::column_view const& column, - 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()); /** @@ -227,7 +227,7 @@ std::unique_ptr days_in_month( */ std::unique_ptr extract_quarter( cudf::column_view const& column, - 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()); /** @@ -258,7 +258,7 @@ enum class rounding_frequency : int32_t { std::unique_ptr ceil_datetimes( cudf::column_view const& column, rounding_frequency freq, - 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()); /** @@ -275,7 +275,7 @@ std::unique_ptr ceil_datetimes( std::unique_ptr floor_datetimes( cudf::column_view const& column, rounding_frequency freq, - 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()); /** @@ -292,7 +292,7 @@ std::unique_ptr floor_datetimes( std::unique_ptr round_datetimes( cudf::column_view const& column, rounding_frequency freq, - 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()); /** @} */ // end of group diff --git a/cpp/include/cudf/detail/aggregation/aggregation.hpp b/cpp/include/cudf/detail/aggregation/aggregation.hpp index 85583a9938d5..ae5078642720 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.hpp +++ b/cpp/include/cudf/detail/aggregation/aggregation.hpp @@ -1382,7 +1382,7 @@ bool is_valid_aggregation(data_type source, aggregation::Kind k); */ void initialize_with_identity(mutable_table_view const& table, host_span aggs, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); } // namespace detail } // namespace cudf diff --git a/cpp/include/cudf/detail/algorithms/copy_if.cuh b/cpp/include/cudf/detail/algorithms/copy_if.cuh index 15455e42637e..43a63516a05d 100644 --- a/cpp/include/cudf/detail/algorithms/copy_if.cuh +++ b/cpp/include/cudf/detail/algorithms/copy_if.cuh @@ -8,7 +8,6 @@ #include #include -#include #include #include @@ -49,7 +48,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); @@ -65,7 +64,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()); @@ -78,7 +77,7 @@ OutputIterator copy_if(InputIterator begin, num_selected.data(), num_items, predicate, - stream.value())); + stream.get())); return result + num_selected.value(stream); } @@ -106,7 +105,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); @@ -123,7 +122,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( @@ -137,7 +136,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); @@ -155,18 +154,18 @@ void copy_if_async(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); auto tmp_bytes = std::size_t{0}; auto no_out = cuda::make_discard_iterator(); CUDF_CUDA_TRY(cub::DeviceSelect::If( - nullptr, tmp_bytes, begin, output, no_out, num_items, predicate, stream.value())); + nullptr, tmp_bytes, begin, output, no_out, num_items, predicate, stream.get())); auto tmp_stg = rmm::device_buffer(tmp_bytes, stream, cudf::get_current_device_resource_ref()); CUDF_CUDA_TRY(cub::DeviceSelect::If( - tmp_stg.data(), tmp_bytes, begin, output, no_out, num_items, predicate, stream.value())); + tmp_stg.data(), tmp_bytes, begin, output, no_out, num_items, predicate, stream.get())); } /** @@ -185,18 +184,18 @@ void copy_if_async(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); auto tmp_bytes = std::size_t{0}; auto no_out = cuda::make_discard_iterator(); CUDF_CUDA_TRY(cub::DeviceSelect::FlaggedIf( - nullptr, tmp_bytes, begin, stencil, result, no_out, num_items, predicate, stream.value())); + nullptr, tmp_bytes, begin, stencil, result, no_out, num_items, predicate, stream.get())); auto tmp = rmm::device_buffer(tmp_bytes, stream, cudf::get_current_device_resource_ref()); CUDF_CUDA_TRY(cub::DeviceSelect::FlaggedIf( - tmp.data(), tmp_bytes, begin, stencil, result, no_out, num_items, predicate, stream.value())); + tmp.data(), tmp_bytes, begin, stencil, result, no_out, num_items, predicate, stream.get())); } } // namespace cudf::detail diff --git a/cpp/include/cudf/detail/algorithms/reduce.cuh b/cpp/include/cudf/detail/algorithms/reduce.cuh index c50062ba10cc..533949254b14 100644 --- a/cpp/include/cudf/detail/algorithms/reduce.cuh +++ b/cpp/include/cudf/detail/algorithms/reduce.cuh @@ -8,7 +8,6 @@ #include #include -#include #include #include @@ -39,11 +38,8 @@ namespace cudf::detail { template > -OutputType reduce(InputIterator begin, - InputIterator end, - OutputType init, - Op binary_op, - rmm::cuda_stream_view stream) +OutputType reduce( + InputIterator begin, InputIterator end, OutputType init, Op binary_op, cuda::stream_ref stream) { auto const num_items = cuda::std::distance(begin, end); @@ -53,7 +49,7 @@ OutputType reduce(InputIterator begin, // Build environment with stream and memory resource for cub::DeviceReduce::Reduce auto env = cuda::std::execution::env{ - cuda::std::execution::prop{cuda::get_stream_t{}, cuda::stream_ref{stream.value()}}, + cuda::std::execution::prop{cuda::get_stream_t{}, cuda::stream_ref{stream.get()}}, cuda::std::execution::prop{cuda::mr::get_memory_resource_t{}, cudf::get_current_device_resource_ref()}}; CUDF_CUDA_TRY(cub::DeviceReduce::Reduce(begin, result.data(), num_items, binary_op, init, env)); @@ -98,7 +94,7 @@ cuda::std::pair reduce_by_key( KeysOutputIterator keys_output, ValuesOutputIterator values_output, Op op, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const num_items = cuda::std::distance(keys_begin, keys_end); @@ -117,7 +113,7 @@ cuda::std::pair reduce_by_key( d_num_runs.data(), op, num_items, - stream.value())); + stream.get())); // Allocate temporary storage rmm::device_buffer d_temp_storage( @@ -133,7 +129,7 @@ cuda::std::pair reduce_by_key( d_num_runs.data(), op, num_items, - stream.value())); + stream.get())); // Copy number of runs back to host via pinned memory auto const num_runs = d_num_runs.value(stream); @@ -160,7 +156,7 @@ void reduce_by_key_async(KeysInputIterator keys_begin, KeysOutputIterator keys_output, ValuesOutputIterator values_output, Op op, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const num_items = cuda::std::distance(keys_begin, keys_end); @@ -174,7 +170,7 @@ void reduce_by_key_async(KeysInputIterator keys_begin, cuda::make_discard_iterator(), op, num_items, - stream.value())); + stream.get())); rmm::device_buffer d_temp_storage( temp_storage_bytes, stream, cudf::get_current_device_resource_ref()); @@ -188,7 +184,7 @@ void reduce_by_key_async(KeysInputIterator keys_begin, cuda::make_discard_iterator(), op, num_items, - stream.value())); + stream.get())); } /** @@ -221,7 +217,7 @@ OutputType transform_reduce(InputIterator begin, TransformationOp transform_op, OutputType init, ReductionOp reduce_op, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const num_items = cuda::std::distance(begin, end); @@ -238,7 +234,7 @@ OutputType transform_reduce(InputIterator begin, reduce_op, transform_op, init, - stream.value())); + stream.get())); rmm::device_buffer d_temp_storage( temp_storage_bytes, stream, cudf::get_current_device_resource_ref()); @@ -250,7 +246,7 @@ OutputType transform_reduce(InputIterator begin, reduce_op, transform_op, init, - stream.value())); + stream.get())); // Copy result back to host via pinned memory return result.value(stream); @@ -272,7 +268,7 @@ OutputType transform_reduce(InputIterator begin, * @return true if the predicate is true for all elements, false otherwise */ template -bool all_of(InputIterator begin, InputIterator end, TransformOp op, rmm::cuda_stream_view stream) +bool all_of(InputIterator begin, InputIterator end, TransformOp op, cuda::stream_ref stream) { return transform_reduce(begin, end, op, true, cuda::std::logical_and{}, stream); } @@ -293,7 +289,7 @@ bool all_of(InputIterator begin, InputIterator end, TransformOp op, rmm::cuda_st * @return true if the predicate is true for any element, false otherwise */ template -bool any_of(InputIterator begin, InputIterator end, TransformOp op, rmm::cuda_stream_view stream) +bool any_of(InputIterator begin, InputIterator end, TransformOp op, cuda::stream_ref stream) { return transform_reduce(begin, end, op, false, cuda::std::logical_or{}, stream); } @@ -314,7 +310,7 @@ bool any_of(InputIterator begin, InputIterator end, TransformOp op, rmm::cuda_st * @return true if the predicate is false for all elements, false otherwise */ template -bool none_of(InputIterator begin, InputIterator end, TransformOp op, rmm::cuda_stream_view stream) +bool none_of(InputIterator begin, InputIterator end, TransformOp op, cuda::stream_ref stream) { return not any_of(begin, end, op, stream); } @@ -339,7 +335,7 @@ template cuda::std::size_t count_if(InputIterator begin, InputIterator end, Predicate predicate, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // Transform each element to 0 or 1 based on predicate, then sum auto transform_op = [predicate] __device__(auto const& val) -> cuda::std::size_t { diff --git a/cpp/include/cudf/detail/calendrical_month_sequence.cuh b/cpp/include/cudf/detail/calendrical_month_sequence.cuh index 46004ca92efe..9543487853db 100644 --- a/cpp/include/cudf/detail/calendrical_month_sequence.cuh +++ b/cpp/include/cudf/detail/calendrical_month_sequence.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -13,10 +13,10 @@ #include #include -#include #include #include +#include #include namespace cudf { @@ -26,7 +26,7 @@ struct calendrical_month_sequence_functor { std::unique_ptr operator()(size_type n, scalar const& input, size_type months, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(cudf::is_timestamp_t::value) { diff --git a/cpp/include/cudf/detail/copy_if.cuh b/cpp/include/cudf/detail/copy_if.cuh index 7f2dfb0e4817..fcf8cd4bb431 100644 --- a/cpp/include/cudf/detail/copy_if.cuh +++ b/cpp/include/cudf/detail/copy_if.cuh @@ -14,12 +14,12 @@ #include #include -#include #include #include #include #include +#include #include namespace cudf { @@ -43,7 +43,7 @@ namespace detail { template std::unique_ptr copy_if(table_view const& input, Filter filter, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/include/cudf/detail/copy_if_else.cuh b/cpp/include/cudf/detail/copy_if_else.cuh index a2b5e37d2715..0462705fbca7 100644 --- a/cpp/include/cudf/detail/copy_if_else.cuh +++ b/cpp/include/cudf/detail/copy_if_else.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -140,7 +140,7 @@ std::unique_ptr copy_if_else(bool nullable, RightIter rhs, FilterFn filter, cudf::data_type output_type, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // This is the type of the cuda::std::optional element in the passed iterators @@ -163,14 +163,14 @@ std::unique_ptr copy_if_else(bool nullable, // call the kernel copy_if_else_kernel - <<>>( + <<>>( lhs_begin, rhs, filter, *out_v, valid_count.data()); out->set_null_count(size - valid_count.value(stream)); } else { // call the kernel copy_if_else_kernel - <<>>(lhs_begin, rhs, filter, *out_v, nullptr); + <<>>(lhs_begin, rhs, filter, *out_v, nullptr); } return out; diff --git a/cpp/include/cudf/detail/datetime.hpp b/cpp/include/cudf/detail/datetime.hpp index 9044f0c3abdb..4557e24740c5 100644 --- a/cpp/include/cudf/detail/datetime.hpp +++ b/cpp/include/cudf/detail/datetime.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -17,63 +17,63 @@ namespace datetime { namespace detail { /** * @copydoc cudf::extract_datetime_component(cudf::column_view const&, datetime_component, - * rmm::cuda_stream_view, rmm::device_async_resource_ref) + * cuda::stream_ref, rmm::device_async_resource_ref) * */ std::unique_ptr extract_datetime_component(cudf::column_view const& column, datetime_component component, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** - * @copydoc cudf::last_day_of_month(cudf::column_view const&, rmm::cuda_stream_view, + * @copydoc cudf::last_day_of_month(cudf::column_view const&, cuda::stream_ref, * rmm::device_async_resource_ref) * */ std::unique_ptr last_day_of_month(cudf::column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** - * @copydoc cudf::day_of_year(cudf::column_view const&, rmm::cuda_stream_view, + * @copydoc cudf::day_of_year(cudf::column_view const&, cuda::stream_ref, * rmm::device_async_resource_ref) * */ std::unique_ptr day_of_year(cudf::column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @copydoc cudf::add_calendrical_months(cudf::column_view const&, cudf::column_view const&, - * rmm::cuda_stream_view, rmm::device_async_resource_ref) + * cuda::stream_ref, rmm::device_async_resource_ref) * */ std::unique_ptr add_calendrical_months(cudf::column_view const& timestamps, cudf::column_view const& months, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @copydoc cudf::add_calendrical_months(cudf::column_view const&, cudf::scalar const&, - * rmm::cuda_stream_view, rmm::device_async_resource_ref) + * cuda::stream_ref, rmm::device_async_resource_ref) * */ std::unique_ptr add_calendrical_months(cudf::column_view const& timestamps, cudf::scalar const& months, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** - * @copydoc cudf::is_leap_year(cudf::column_view const&, rmm::cuda_stream_view, + * @copydoc cudf::is_leap_year(cudf::column_view const&, cuda::stream_ref, * rmm::device_async_resource_ref) * */ std::unique_ptr is_leap_year(cudf::column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr extract_quarter(cudf::column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/detail/device_scalar.hpp b/cpp/include/cudf/detail/device_scalar.hpp index e2380ef8cb1d..26df803549a6 100644 --- a/cpp/include/cudf/detail/device_scalar.hpp +++ b/cpp/include/cudf/detail/device_scalar.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -9,10 +9,11 @@ #include #include -#include #include #include +#include + namespace CUDF_EXPORT cudf { namespace detail { @@ -33,7 +34,7 @@ class device_scalar : public rmm::device_scalar { device_scalar() = delete; explicit device_scalar( - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()) : rmm::device_scalar(stream, mr), bounce_buffer{make_pinned_vector(1, stream)} { @@ -41,7 +42,7 @@ class device_scalar : public rmm::device_scalar { explicit device_scalar( T const& initial_value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()) : rmm::device_scalar(stream, mr), bounce_buffer{make_pinned_vector(1, stream)} { @@ -50,31 +51,31 @@ class device_scalar : public rmm::device_scalar { } device_scalar(device_scalar const& other, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()) : rmm::device_scalar(other, stream, mr), bounce_buffer{make_pinned_vector(1, stream)} { } - [[nodiscard]] T value(rmm::cuda_stream_view stream) const + [[nodiscard]] T value(cuda::stream_ref stream) const { cuda_memcpy(bounce_buffer, device_span(this->data(), 1), stream); return std::move(bounce_buffer[0]); } - void set_value_async(T const& value, rmm::cuda_stream_view stream) + void set_value_async(T const& value, cuda::stream_ref stream) { bounce_buffer[0] = value; cuda_memcpy_async(device_span(this->data(), 1), bounce_buffer, stream); } - void set_value_async(T&& value, rmm::cuda_stream_view stream) + void set_value_async(T&& value, cuda::stream_ref stream) { bounce_buffer[0] = std::move(value); cuda_memcpy_async(device_span{this->data(), 1}, bounce_buffer, stream); } - void set_value_to_zero_async(rmm::cuda_stream_view stream) { set_value_async(T{}, stream); } + void set_value_to_zero_async(cuda::stream_ref stream) { set_value_async(T{}, stream); } private: mutable cudf::detail::host_vector bounce_buffer; diff --git a/cpp/include/cudf/detail/gather.cuh b/cpp/include/cudf/detail/gather.cuh index e5bb1f9ff575..8dddd0090c16 100644 --- a/cpp/include/cudf/detail/gather.cuh +++ b/cpp/include/cudf/detail/gather.cuh @@ -27,10 +27,10 @@ #include #include -#include #include #include +#include #include #include @@ -112,7 +112,7 @@ void gather_helper(InputItr source_itr, MapIterator gather_map_begin, MapIterator gather_map_end, bool nullify_out_of_bounds, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { using map_type = typename std::iterator_traits::value_type; if (nullify_out_of_bounds) { @@ -166,7 +166,7 @@ struct column_gatherer { MapIterator gather_map_begin, MapIterator gather_map_end, bool nullify_out_of_bounds, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { column_gatherer_impl gatherer{}; @@ -206,7 +206,7 @@ struct column_gatherer_impl { MapItType gather_map_begin, MapItType gather_map_end, bool nullify_out_of_bounds, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (true == nullify_out_of_bounds) { @@ -325,7 +325,7 @@ struct column_gatherer_impl { MapItRoot gather_map_begin, MapItRoot gather_map_end, bool nullify_out_of_bounds, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { lists_column_view list(column); @@ -388,7 +388,7 @@ struct column_gatherer_impl { MapItType gather_map_begin, MapItType gather_map_end, bool nullify_out_of_bounds, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { dictionary_column_view dictionary(source_column); @@ -425,7 +425,7 @@ struct column_gatherer_impl { MapItRoot gather_map_begin, MapItRoot gather_map_end, bool nullify_out_of_bounds, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const gather_map_size = std::distance(gather_map_begin, gather_map_end); @@ -510,7 +510,7 @@ void gather_bitmask(table_device_view input, size_type mask_count, size_type mask_size, size_type* valid_counts, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (mask_size == 0) { return; } @@ -522,7 +522,7 @@ void gather_bitmask(table_device_view input, valid_if_n_kernel; cudf::detail::grid_1d grid{mask_size, block_size, 1}; - kernel<<>>( + kernel<<>>( counting_it, counting_it, selector, masks, mask_count, mask_size, valid_counts); } @@ -531,7 +531,7 @@ void gather_bitmask(table_view const& source, MapIterator gather_map, std::vector>& target, gather_bitmask_op op, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (target.empty()) { return; } @@ -629,7 +629,7 @@ std::unique_ptr
gather(table_view const& source_table, MapIterator gather_map_begin, MapIterator gather_map_end, out_of_bounds_policy bounds_policy, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { std::vector> destination_columns; diff --git a/cpp/include/cudf/detail/get_value.cuh b/cpp/include/cudf/detail/get_value.cuh index 7d91b36aebac..1b49b98915df 100644 --- a/cpp/include/cudf/detail/get_value.cuh +++ b/cpp/include/cudf/detail/get_value.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -11,7 +11,7 @@ #include #include -#include +#include namespace cudf { namespace detail { @@ -32,7 +32,7 @@ namespace detail { * @return Value from the `col_view[element_index]` */ template -T get_value(column_view const& col_view, size_type element_index, rmm::cuda_stream_view stream) +T get_value(column_view const& col_view, size_type element_index, cuda::stream_ref stream) { CUDF_EXPECTS(cudf::is_fixed_width(col_view.type()), "get_value supports only fixed-width types"); CUDF_EXPECTS(data_type(type_to_id()) == col_view.type(), "get_value data type mismatch"); diff --git a/cpp/include/cudf/detail/indexalator.cuh b/cpp/include/cudf/detail/indexalator.cuh index 55b0b483c1fd..177debbe3a61 100644 --- a/cpp/include/cudf/detail/indexalator.cuh +++ b/cpp/include/cudf/detail/indexalator.cuh @@ -11,11 +11,11 @@ #include #include -#include - #include #include #include +#include +#include namespace cudf { namespace detail { @@ -318,7 +318,7 @@ struct indexalator_factory { /** * @brief Create an accessor from a scalar. */ - scalar_optional_index_accessor(scalar const& input, rmm::cuda_stream_view stream) + scalar_optional_index_accessor(scalar const& input, cuda::stream_ref stream) : is_null{!input.is_valid(stream)} { iter = indexalator_factory::make_input_iterator(input); @@ -341,7 +341,7 @@ struct indexalator_factory { /** * @brief Create an index iterator with an optional index accessor for a scalar. */ - static auto make_input_optional_iterator(scalar const& input, rmm::cuda_stream_view stream) + static auto make_input_optional_iterator(scalar const& input, cuda::stream_ref stream) { return cuda::transform_iterator(cuda::make_constant_iterator(0), scalar_optional_index_accessor{input, stream}); diff --git a/cpp/include/cudf/detail/interop.hpp b/cpp/include/cudf/detail/interop.hpp index 65a12102eae3..2e857aa054a7 100644 --- a/cpp/include/cudf/detail/interop.hpp +++ b/cpp/include/cudf/detail/interop.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -9,7 +9,7 @@ #include #include -#include +#include namespace cudf { namespace detail { @@ -18,14 +18,14 @@ namespace detail { * @copydoc cudf::from_dlpack */ std::unique_ptr
from_dlpack(DLManagedTensor const* managed_tensor, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @copydoc cudf::to_dlpack */ DLManagedTensor* to_dlpack(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/detail/label_bins.hpp b/cpp/include/cudf/detail/label_bins.hpp index 6f63bfa15a59..23e1fbca598e 100644 --- a/cpp/include/cudf/detail/label_bins.hpp +++ b/cpp/include/cudf/detail/label_bins.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -12,7 +12,7 @@ #include #include -#include +#include namespace CUDF_EXPORT cudf { @@ -27,7 +27,7 @@ namespace detail { /** * @copydoc cudf::label_bins(column_view const& input, column_view const& left_edges, inclusive - * left_inclusive, column_view const& right_edges, inclusive right_inclusive, rmm::cuda_stream_view, + * left_inclusive, column_view const& right_edges, inclusive right_inclusive, cuda::stream_ref, * rmm::device_async_resource_ref mr) * * @param stream Stream view on which to allocate resources and queue execution. @@ -37,7 +37,7 @@ std::unique_ptr label_bins(column_view const& input, inclusive left_inclusive, column_view const& right_edges, inclusive right_inclusive, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @} */ // end of group diff --git a/cpp/include/cudf/detail/labeling/label_segments.cuh b/cpp/include/cudf/detail/labeling/label_segments.cuh index b379117273b0..d9ff008e5237 100644 --- a/cpp/include/cudf/detail/labeling/label_segments.cuh +++ b/cpp/include/cudf/detail/labeling/label_segments.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -9,11 +9,11 @@ #include #include -#include #include #include #include +#include #include #include #include @@ -57,7 +57,7 @@ void label_segments(InputIterator offsets_begin, InputIterator offsets_end, OutputIterator label_begin, OutputIterator label_end, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const num_labels = cuda::std::distance(label_begin, label_end); @@ -139,7 +139,7 @@ void labels_to_offsets(InputIterator labels_begin, InputIterator labels_end, OutputIterator offsets_begin, OutputIterator offsets_end, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // Always fill the entire output array with `0` value regardless of the input. using OutputType = cuda::std::iter_value_t; diff --git a/cpp/include/cudf/detail/merge.hpp b/cpp/include/cudf/detail/merge.hpp index d2451c8ae048..130b209aea4d 100644 --- a/cpp/include/cudf/detail/merge.hpp +++ b/cpp/include/cudf/detail/merge.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -36,7 +36,7 @@ using index_vector = rmm::device_uvector; * std::vector const& key_cols, * std::vector const& column_order, * std::vector const& null_precedence, - * rmm::cuda_stream_view stream, + * cuda::stream_ref stream, * rmm::device_async_resource_ref mr) * * @param stream CUDA stream used for device memory operations and kernel launches @@ -45,7 +45,7 @@ std::unique_ptr merge(std::vector const& tables_to_merg std::vector const& key_cols, std::vector const& column_order, std::vector const& null_precedence, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/detail/quantiles.hpp b/cpp/include/cudf/detail/quantiles.hpp index 6a73ffe6f5c1..bc7f62e5a7e2 100644 --- a/cpp/include/cudf/detail/quantiles.hpp +++ b/cpp/include/cudf/detail/quantiles.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -9,7 +9,7 @@ #include #include -#include +#include namespace cudf { namespace detail { @@ -22,7 +22,7 @@ std::unique_ptr quantile(column_view const& input, interpolation interp, column_view const& ordered_indices, bool exact, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -34,7 +34,7 @@ std::unique_ptr
quantiles(table_view const& input, cudf::sorted is_input_sorted, std::vector const& column_order, std::vector const& null_precedence, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -43,7 +43,7 @@ std::unique_ptr
quantiles(table_view const& input, */ std::unique_ptr percentile_approx(tdigest::tdigest_column_view const& input, column_view const& percentiles, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/detail/replace.hpp b/cpp/include/cudf/detail/replace.hpp index 49a9a6daa7d1..06f2bcb0e131 100644 --- a/cpp/include/cudf/detail/replace.hpp +++ b/cpp/include/cudf/detail/replace.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -9,7 +9,7 @@ #include #include -#include +#include #include @@ -17,47 +17,47 @@ namespace cudf { namespace detail { /** * @copydoc cudf::replace_nulls(column_view const&, column_view const&, - * rmm::cuda_stream_view, rmm::device_async_resource_ref) + * cuda::stream_ref, rmm::device_async_resource_ref) */ std::unique_ptr replace_nulls(column_view const& input, cudf::column_view const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @copydoc cudf::replace_nulls(column_view const&, scalar const&, - * rmm::cuda_stream_view, rmm::device_async_resource_ref) + * cuda::stream_ref, rmm::device_async_resource_ref) */ std::unique_ptr replace_nulls(column_view const& input, scalar const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @copydoc cudf::replace_nulls(column_view const&, replace_policy const&, - * rmm::cuda_stream_view, rmm::device_async_resource_ref) + * cuda::stream_ref, rmm::device_async_resource_ref) */ std::unique_ptr replace_nulls(column_view const& input, replace_policy const& replace_policy, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @copydoc cudf::replace_nans(column_view const&, column_view const&, - * rmm::cuda_stream_view, rmm::device_async_resource_ref) + * cuda::stream_ref, rmm::device_async_resource_ref) */ std::unique_ptr replace_nans(column_view const& input, column_view const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @copydoc cudf::replace_nans(column_view const&, scalar const&, - * rmm::cuda_stream_view, rmm::device_async_resource_ref) + * cuda::stream_ref, rmm::device_async_resource_ref) */ std::unique_ptr replace_nans(column_view const& input, scalar const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -66,14 +66,14 @@ std::unique_ptr replace_nans(column_view const& input, std::unique_ptr find_and_replace_all(column_view const& input_col, column_view const& values_to_replace, column_view const& replacement_values, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @copydoc cudf::normalize_nans_and_zeros */ std::unique_ptr normalize_nans_and_zeros(column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/detail/reshape.hpp b/cpp/include/cudf/detail/reshape.hpp index b114a5b4501b..54e04ca82cc7 100644 --- a/cpp/include/cudf/detail/reshape.hpp +++ b/cpp/include/cudf/detail/reshape.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -10,7 +10,7 @@ #include #include -#include +#include #include @@ -21,14 +21,14 @@ namespace detail { */ std::unique_ptr
tile(table_view const& input, size_type count, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref mr); /** * @copydoc cudf::interleave_columns */ std::unique_ptr interleave_columns(table_view const& input, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref mr); /** @@ -36,7 +36,7 @@ std::unique_ptr interleave_columns(table_view const& input, */ void table_to_array(table_view const& input, device_span output, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); } // namespace detail } // namespace cudf diff --git a/cpp/include/cudf/detail/row_operator/equality.cuh b/cpp/include/cudf/detail/row_operator/equality.cuh index 9a5fd68b6ceb..9ab9aa260a01 100644 --- a/cpp/include/cudf/detail/row_operator/equality.cuh +++ b/cpp/include/cudf/detail/row_operator/equality.cuh @@ -410,7 +410,7 @@ class self_comparator { * @param temp_mr Device memory resource used for temporary allocations */ self_comparator(table_view const& t, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) : d_t(preprocessed_table::create(t, stream, temp_mr)) { @@ -522,7 +522,7 @@ class two_table_comparator { */ two_table_comparator(table_view const& left, table_view const& right, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr); /** diff --git a/cpp/include/cudf/detail/row_operator/hashing.cuh b/cpp/include/cudf/detail/row_operator/hashing.cuh index 5cd574eb4342..a98c5d591274 100644 --- a/cpp/include/cudf/detail/row_operator/hashing.cuh +++ b/cpp/include/cudf/detail/row_operator/hashing.cuh @@ -242,9 +242,7 @@ class row_hasher { * comparisons using this object. * @param temp_mr Device memory resource used for temporary allocations */ - row_hasher(table_view const& t, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref temp_mr) + row_hasher(table_view const& t, cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) : d_t(preprocessed_table::create(t, stream, temp_mr)) { } diff --git a/cpp/include/cudf/detail/row_operator/lexicographic.cuh b/cpp/include/cudf/detail/row_operator/lexicographic.cuh index 8882b2164d0f..395464f21274 100644 --- a/cpp/include/cudf/detail/row_operator/lexicographic.cuh +++ b/cpp/include/cudf/detail/row_operator/lexicographic.cuh @@ -679,7 +679,7 @@ struct less_equivalent_comparator struct preprocessed_table { using table_device_view_owner = std::invoke_result_t; /** @@ -707,7 +707,7 @@ struct preprocessed_table { static std::shared_ptr create(table_view const& table, host_span column_order, host_span null_precedence, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @brief Preprocess tables for use with lexicographical comparison @@ -735,7 +735,7 @@ struct preprocessed_table { table_view const& rhs, host_span column_order, host_span null_precedence, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); private: friend class self_comparator; @@ -767,7 +767,7 @@ struct preprocessed_table { host_span column_order, host_span null_precedence, bool has_ranked_children, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @brief Construct a preprocessed table for use with lexicographical comparison @@ -924,7 +924,7 @@ class self_comparator { self_comparator(table_view const& t, host_span column_order = {}, host_span null_precedence = {}, - rmm::cuda_stream_view stream = cudf::get_default_stream()) + cuda::stream_ref stream = cudf::get_default_stream()) : d_t{preprocessed_table::create(t, column_order, null_precedence, stream)} { } @@ -1081,7 +1081,7 @@ class two_table_comparator { table_view const& right, host_span column_order = {}, host_span null_precedence = {}, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Construct an owning object for performing a lexicographic comparison between two rows of diff --git a/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh b/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh index 08949be1f0b7..ba85a2b0578a 100644 --- a/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh +++ b/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh @@ -7,11 +7,12 @@ #include -#include #include #include #include +#include + #include #include @@ -52,7 +53,7 @@ struct preprocessed_table { * @return A preprocessed table as shared pointer */ static std::shared_ptr create(table_view const& table, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr); /** @@ -73,7 +74,7 @@ struct preprocessed_table { using table_device_view_owner = std::invoke_result_t; preprocessed_table(table_device_view_owner&& table, diff --git a/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh b/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh index 79f50d9339dc..4d3feb6f1ccd 100644 --- a/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh +++ b/cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh @@ -11,11 +11,11 @@ #include #include -#include #include #include #include +#include #include #include @@ -256,7 +256,7 @@ auto sizes_to_offsets(SizesIterator begin, SizesIterator end, OffsetsIterator result, int64_t initial_offset, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { using SizeType = cuda::std::iter_value_t; static_assert(std::is_integral_v, @@ -302,7 +302,7 @@ template std::pair, size_type> make_offsets_child_column( InputIterator begin, InputIterator end, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto count = static_cast(std::distance(begin, end)); diff --git a/cpp/include/cudf/detail/stream_compaction.hpp b/cpp/include/cudf/detail/stream_compaction.hpp index e506d34a5cc1..618594904713 100644 --- a/cpp/include/cudf/detail/stream_compaction.hpp +++ b/cpp/include/cudf/detail/stream_compaction.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -11,29 +11,30 @@ #include #include -#include #include +#include + namespace cudf { namespace detail { /** * @copydoc cudf::drop_nulls(table_view const&, std::vector const&, - * cudf::size_type, rmm::cuda_stream_view, rmm::device_async_resource_ref) + * cudf::size_type, cuda::stream_ref, rmm::device_async_resource_ref) */ std::unique_ptr
drop_nulls(table_view const& input, std::vector const& keys, cudf::size_type keep_threshold, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @copydoc cudf::drop_nans(table_view const&, std::vector const&, - * cudf::size_type, rmm::cuda_stream_view, rmm::device_async_resource_ref) + * cudf::size_type, cuda::stream_ref, rmm::device_async_resource_ref) */ std::unique_ptr
drop_nans(table_view const& input, std::vector const& keys, cudf::size_type keep_threshold, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -52,7 +53,7 @@ enum class mask_type : bool { std::unique_ptr
apply_mask(table_view const& input, column_view const& boolean_mask, mask_type mask_kind, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -62,7 +63,7 @@ std::unique_ptr
unique(table_view const& input, std::vector const& keys, duplicate_keep_option keep, null_equality nulls_equal, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -73,7 +74,7 @@ std::unique_ptr
distinct(table_view const& input, duplicate_keep_option keep, null_equality nulls_equal, nan_equality nans_equal, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -84,7 +85,7 @@ std::unique_ptr
stable_distinct(table_view const& input, duplicate_keep_option keep, null_equality nulls_equal, nan_equality nans_equal, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -96,7 +97,7 @@ rmm::device_uvector distinct_indices(table_view const& input, duplicate_keep_option keep, null_equality nulls_equal, nan_equality nans_equal, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/detail/tdigest/tdigest.hpp b/cpp/include/cudf/detail/tdigest/tdigest.hpp index 9e2fe5ee8a08..3485c9cfcf4b 100644 --- a/cpp/include/cudf/detail/tdigest/tdigest.hpp +++ b/cpp/include/cudf/detail/tdigest/tdigest.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -11,7 +11,7 @@ #include #include -#include +#include namespace CUDF_EXPORT cudf { namespace tdigest::detail { @@ -62,7 +62,7 @@ std::unique_ptr group_tdigest(column_view const& values, cudf::device_span group_valid_counts, size_type num_groups, int max_centroids, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -105,7 +105,7 @@ std::unique_ptr group_merge_tdigest(column_view const& values, cudf::device_span group_labels, size_type num_groups, int max_centroids, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -131,7 +131,7 @@ std::unique_ptr make_tdigest_column(size_type num_rows, std::unique_ptr&& tdigest_offsets, std::unique_ptr&& min_values, std::unique_ptr&& max_values, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -147,7 +147,7 @@ std::unique_ptr make_tdigest_column(size_type num_rows, */ CUDF_EXPORT std::unique_ptr make_empty_tdigests_column(size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -160,7 +160,7 @@ std::unique_ptr make_empty_tdigests_column(size_type num_rows, * * @returns A scalar of an empty tdigest cluster. */ -std::unique_ptr make_empty_tdigest_scalar(rmm::cuda_stream_view stream, +std::unique_ptr make_empty_tdigest_scalar(cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -194,7 +194,7 @@ std::unique_ptr make_empty_tdigest_scalar(rmm::cuda_stream_view stream, */ std::unique_ptr reduce_tdigest(column_view const& values, int max_centroids, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -228,7 +228,7 @@ std::unique_ptr reduce_tdigest(column_view const& values, */ std::unique_ptr reduce_merge_tdigest(column_view const& input, int max_centroids, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace tdigest::detail diff --git a/cpp/include/cudf/detail/timezone.hpp b/cpp/include/cudf/detail/timezone.hpp index a5352c189788..5f7d32f7fe89 100644 --- a/cpp/include/cudf/detail/timezone.hpp +++ b/cpp/include/cudf/detail/timezone.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -9,20 +9,20 @@ #include #include -#include +#include namespace cudf { namespace detail { /** * @copydoc cudf::make_timezone_transition_table(std::optional, std::string_view, - * rmm::cuda_stream_view, rmm::device_async_resource_ref) + * cuda::stream_ref, rmm::device_async_resource_ref) * */ std::unique_ptr
make_timezone_transition_table( std::optional tzif_dir, std::string_view timezone_name, - 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 detail diff --git a/cpp/include/cudf/detail/transpose.hpp b/cpp/include/cudf/detail/transpose.hpp index bfa13ab74780..f9d37901e6e4 100644 --- a/cpp/include/cudf/detail/transpose.hpp +++ b/cpp/include/cudf/detail/transpose.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -9,7 +9,7 @@ #include #include -#include +#include namespace cudf { namespace detail { @@ -17,7 +17,7 @@ namespace detail { * @copydoc cudf::transpose */ std::pair, table_view> transpose(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/detail/utilities/batched_memcpy.hpp b/cpp/include/cudf/detail/utilities/batched_memcpy.hpp index 4b0ba06c6991..6b158fe0b8b1 100644 --- a/cpp/include/cudf/detail/utilities/batched_memcpy.hpp +++ b/cpp/include/cudf/detail/utilities/batched_memcpy.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -8,11 +8,11 @@ #include #include -#include #include #include #include +#include namespace CUDF_EXPORT cudf { namespace detail { @@ -36,13 +36,13 @@ void batched_memcpy_async(SrcIterator src_iter, DstIterator dst_iter, SizeIterator size_iter, size_t num_buffs, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { size_t temp_storage_bytes = 0; cub::DeviceMemcpy::Batched( - nullptr, temp_storage_bytes, src_iter, dst_iter, size_iter, num_buffs, stream.value()); + nullptr, temp_storage_bytes, src_iter, dst_iter, size_iter, num_buffs, stream.get()); - rmm::device_buffer d_temp_storage{temp_storage_bytes, stream.value()}; + rmm::device_buffer d_temp_storage{temp_storage_bytes, stream.get()}; cub::DeviceMemcpy::Batched(d_temp_storage.data(), temp_storage_bytes, @@ -50,7 +50,7 @@ void batched_memcpy_async(SrcIterator src_iter, dst_iter, size_iter, num_buffs, - stream.value()); + stream.get()); } } // namespace detail diff --git a/cpp/include/cudf/detail/utilities/batched_memset.hpp b/cpp/include/cudf/detail/utilities/batched_memset.hpp index 66b3fa7b1df7..a59792181045 100644 --- a/cpp/include/cudf/detail/utilities/batched_memset.hpp +++ b/cpp/include/cudf/detail/utilities/batched_memset.hpp @@ -10,12 +10,13 @@ #include #include -#include #include #include #include #include +#include +#include #include namespace CUDF_EXPORT cudf { @@ -33,7 +34,7 @@ namespace detail { template void batched_memset(cudf::host_span const> host_buffers, T const value, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); @@ -59,14 +60,14 @@ void batched_memset(cudf::host_span const> host_buffers, auto const num_buffers = host_buffers.size(); cub::DeviceCopy::Batched( - nullptr, temp_storage_bytes, iter_in, iter_out, sizes, num_buffers, stream); + nullptr, temp_storage_bytes, iter_in, iter_out, sizes, num_buffers, stream.get()); // Allocate temporary storage rmm::device_buffer d_temp_storage( temp_storage_bytes, stream, cudf::get_current_device_resource_ref()); cub::DeviceCopy::Batched( - d_temp_storage.data(), temp_storage_bytes, iter_in, iter_out, sizes, num_buffers, stream); + d_temp_storage.data(), temp_storage_bytes, iter_in, iter_out, sizes, num_buffers, stream.get()); } } // namespace detail diff --git a/cpp/include/cudf/detail/utilities/cuda.cuh b/cpp/include/cudf/detail/utilities/cuda.cuh index 6dcf879017da..f38975008d69 100644 --- a/cpp/include/cudf/detail/utilities/cuda.cuh +++ b/cpp/include/cudf/detail/utilities/cuda.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -11,10 +11,9 @@ #include -#include - #include #include +#include #endif @@ -88,9 +87,9 @@ CUDF_KERNEL void single_thread_kernel(F f) * @param stream CUDA stream used for the kernel launch */ template -void device_single_thread(Functor functor, rmm::cuda_stream_view stream) +void device_single_thread(Functor functor, cuda::stream_ref stream) { - single_thread_kernel<<<1, 1, 0, stream.value()>>>(functor); + single_thread_kernel<<<1, 1, 0, stream.get()>>>(functor); } #endif diff --git a/cpp/include/cudf/detail/utilities/cuda.hpp b/cpp/include/cudf/detail/utilities/cuda.hpp index 492024f43d74..912d17273149 100644 --- a/cpp/include/cudf/detail/utilities/cuda.hpp +++ b/cpp/include/cudf/detail/utilities/cuda.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -9,6 +9,8 @@ #include #include +#include + #include namespace CUDF_EXPORT cudf { @@ -50,5 +52,17 @@ cudf::size_type elements_per_thread(Kernel kernel, return std::clamp(per_thread, 1, max_per_thread); } +/** + * @brief Synchronize a CUDA stream. + * + * `cuda::stream_ref::sync()` uses the CUDA driver API, which requires the stream's CUDA context to + * be current on the calling thread. Use the runtime API for host-facing synchronization points that + * may be called from threads without a current CUDA context. + */ +inline void sync_stream(cuda::stream_ref stream) +{ + CUDF_CUDA_TRY(cudaStreamSynchronize(stream.get())); +} + } // namespace detail } // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/detail/utilities/cuda_memcpy.hpp b/cpp/include/cudf/detail/utilities/cuda_memcpy.hpp index 6c4997f5504b..1c98c6a235ce 100644 --- a/cpp/include/cudf/detail/utilities/cuda_memcpy.hpp +++ b/cpp/include/cudf/detail/utilities/cuda_memcpy.hpp @@ -5,11 +5,12 @@ #pragma once +#include #include #include #include -#include +#include namespace CUDF_EXPORT cudf { namespace detail { @@ -17,7 +18,7 @@ namespace detail { enum class host_memory_kind : uint8_t { PINNED, PAGEABLE }; void cuda_memcpy_async_impl( - void* dst, void const* src, size_t size, host_memory_kind kind, rmm::cuda_stream_view stream); + void* dst, void const* src, size_t size, host_memory_kind kind, cuda::stream_ref stream); /** * @brief Wrapper around cudaMemcpyBatchAsync @@ -44,7 +45,7 @@ void cuda_memcpy_async_impl( void const* const* srcs, std::size_t const* sizes, std::size_t count, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @brief Asynchronously copies a single buffer, wrapping `memcpy_batch_async`. @@ -65,7 +66,7 @@ void cuda_memcpy_async_impl( [[nodiscard]] cudaError_t memcpy_async(void* dst, void const* src, size_t count, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @brief Asynchronously copies data from host to device memory. @@ -77,7 +78,7 @@ void cuda_memcpy_async_impl( * @param stream CUDA stream used for the copy */ template -void cuda_memcpy_async(device_span dst, host_span src, rmm::cuda_stream_view stream) +void cuda_memcpy_async(device_span dst, host_span src, cuda::stream_ref stream) { CUDF_EXPECTS(dst.size() == src.size(), "Mismatched sizes in cuda_memcpy_async"); auto const is_pinned = src.is_device_accessible(); @@ -98,7 +99,7 @@ void cuda_memcpy_async(device_span dst, host_span src, rmm::cuda_str * @param stream CUDA stream used for the copy */ template -void cuda_memcpy_async(host_span dst, device_span src, rmm::cuda_stream_view stream) +void cuda_memcpy_async(host_span dst, device_span src, cuda::stream_ref stream) { CUDF_EXPECTS(dst.size() == src.size(), "Mismatched sizes in cuda_memcpy_async"); auto const is_pinned = dst.is_device_accessible(); @@ -119,10 +120,10 @@ void cuda_memcpy_async(host_span dst, device_span src, rmm::cuda_str * @param stream CUDA stream used for the copy */ template -void cuda_memcpy(device_span dst, host_span src, rmm::cuda_stream_view stream) +void cuda_memcpy(device_span dst, host_span src, cuda::stream_ref stream) { cuda_memcpy_async(dst, src, stream); - stream.synchronize(); + cudf::detail::sync_stream(stream); } /** @@ -135,10 +136,10 @@ void cuda_memcpy(device_span dst, host_span src, rmm::cuda_stream_vi * @param stream CUDA stream used for the copy */ template -void cuda_memcpy(host_span dst, device_span src, rmm::cuda_stream_view stream) +void cuda_memcpy(host_span dst, device_span src, cuda::stream_ref stream) { cuda_memcpy_async(dst, src, stream); - stream.synchronize(); + cudf::detail::sync_stream(stream); } } // namespace detail diff --git a/cpp/include/cudf/detail/utilities/host_memory.hpp b/cpp/include/cudf/detail/utilities/host_memory.hpp index d12b7c71226e..e4548d253638 100644 --- a/cpp/include/cudf/detail/utilities/host_memory.hpp +++ b/cpp/include/cudf/detail/utilities/host_memory.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -28,7 +28,7 @@ CUDF_EXPORT rmm::host_async_resource_ref get_pageable_memory_resource(); * @return The allocator to be used for the host memory allocation */ template -rmm_host_allocator get_host_allocator(std::size_t size, rmm::cuda_stream_view stream) +rmm_host_allocator get_host_allocator(std::size_t size, cuda::stream_ref stream) { if (size * sizeof(T) <= get_allocate_host_as_pinned_threshold()) { return {get_pinned_memory_resource(), stream}; diff --git a/cpp/include/cudf/detail/utilities/host_vector.hpp b/cpp/include/cudf/detail/utilities/host_vector.hpp index a65a1879ac5a..862604c5430e 100644 --- a/cpp/include/cudf/detail/utilities/host_vector.hpp +++ b/cpp/include/cudf/detail/utilities/host_vector.hpp @@ -5,6 +5,7 @@ #pragma once +#include #include #include #include @@ -100,7 +101,7 @@ class rmm_host_allocator { #pragma nv_exec_check_disable #endif template - rmm_host_allocator(ResourceType _mr, rmm::cuda_stream_view _stream) + rmm_host_allocator(ResourceType _mr, cuda::stream_ref _stream) : mr(std::move(_mr)), stream(std::move(_stream)), _is_device_accessible{ @@ -143,7 +144,7 @@ class rmm_host_allocator { auto const result = mr.allocate(stream, cnt * sizeof(value_type), alignof(value_type)); // Synchronize to ensure the memory is allocated before thrust::host_vector initialization // TODO: replace thrust::host_vector with a type that does not require synchronization - stream.synchronize(); + cudf::detail::sync_stream(stream); return static_cast(result); } @@ -199,7 +200,7 @@ class rmm_host_allocator { private: rmm::host_async_resource_ref mr; - rmm::cuda_stream_view stream; + cuda::stream_ref stream; bool _is_device_accessible; }; diff --git a/cpp/include/cudf/detail/utilities/vector_factories.hpp b/cpp/include/cudf/detail/utilities/vector_factories.hpp index 00c5c4048128..9d1d26d64025 100644 --- a/cpp/include/cudf/detail/utilities/vector_factories.hpp +++ b/cpp/include/cudf/detail/utilities/vector_factories.hpp @@ -10,6 +10,7 @@ * @file vector_factories.hpp */ +#include #include #include #include @@ -20,10 +21,11 @@ #include #include -#include #include #include +#include + #include #include @@ -43,11 +45,11 @@ namespace detail { */ template rmm::device_uvector make_zeroed_device_uvector_async(std::size_t size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { rmm::device_uvector ret(size, stream, mr); - CUDF_CUDA_TRY(cudaMemsetAsync(ret.data(), 0, size * sizeof(T), stream.value())); + CUDF_CUDA_TRY(cudaMemsetAsync(ret.data(), 0, size * sizeof(T), stream.get())); return ret; } @@ -64,12 +66,12 @@ rmm::device_uvector make_zeroed_device_uvector_async(std::size_t size, */ template rmm::device_uvector make_zeroed_device_uvector(std::size_t size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { rmm::device_uvector ret(size, stream, mr); - CUDF_CUDA_TRY(cudaMemsetAsync(ret.data(), 0, size * sizeof(T), stream.value())); - stream.synchronize(); + CUDF_CUDA_TRY(cudaMemsetAsync(ret.data(), 0, size * sizeof(T), stream.get())); + cudf::detail::sync_stream(stream); return ret; } @@ -87,7 +89,7 @@ rmm::device_uvector make_zeroed_device_uvector(std::size_t size, */ template rmm::device_uvector> make_device_uvector_async( - host_span source_data, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + host_span source_data, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { using value_type = std::remove_cv_t; rmm::device_uvector ret(source_data.size(), stream, mr); @@ -109,7 +111,7 @@ rmm::device_uvector> make_device_uvector_async( */ template rmm::device_uvector make_device_uvector_async( - Container const& c, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + Container const& c, cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(std::is_convertible_v>) { return make_device_uvector_async(host_span{c}, stream, mr); @@ -129,7 +131,7 @@ rmm::device_uvector make_device_uvector_async( */ template rmm::device_uvector make_device_uvector_async(std::vector const& source_data, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return make_device_uvector_async(host_span{source_data}, stream, mr); @@ -149,7 +151,7 @@ rmm::device_uvector make_device_uvector_async(std::vector const */ template rmm::device_uvector> make_device_uvector_async( - device_span source_data, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + device_span source_data, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { using value_type = std::remove_cv_t; rmm::device_uvector ret(source_data.size(), stream, mr); @@ -172,7 +174,7 @@ rmm::device_uvector> make_device_uvector_async( */ template rmm::device_uvector make_device_uvector_async( - Container const& c, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + Container const& c, cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(std::is_convertible_v>) { return make_device_uvector_async( @@ -193,11 +195,11 @@ rmm::device_uvector make_device_uvector_async( */ template rmm::device_uvector make_device_uvector(host_span source_data, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto ret = make_device_uvector_async(source_data, stream, mr); - stream.synchronize(); + cudf::detail::sync_stream(stream); return ret; } @@ -215,7 +217,7 @@ rmm::device_uvector make_device_uvector(host_span source_data, */ template rmm::device_uvector make_device_uvector( - Container const& c, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + Container const& c, cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(std::is_convertible_v>) { return make_device_uvector(host_span{c}, stream, mr); @@ -235,7 +237,7 @@ rmm::device_uvector make_device_uvector( */ template rmm::device_uvector make_device_uvector(std::vector const& source_data, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return make_device_uvector(host_span{source_data}, stream, mr); @@ -255,11 +257,11 @@ rmm::device_uvector make_device_uvector(std::vector const& sour */ template rmm::device_uvector make_device_uvector(device_span source_data, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto ret = make_device_uvector_async(source_data, stream, mr); - stream.synchronize(); + cudf::detail::sync_stream(stream); return ret; } @@ -277,7 +279,7 @@ rmm::device_uvector make_device_uvector(device_span source_data, */ template rmm::device_uvector make_device_uvector( - Container const& c, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + Container const& c, cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(std::is_convertible_v>) { return make_device_uvector(device_span{c}, stream, mr); @@ -295,8 +297,7 @@ rmm::device_uvector make_device_uvector( * @return The data copied to the host */ template -std::vector> make_std_vector_async(device_span v, - rmm::cuda_stream_view stream) +std::vector> make_std_vector_async(device_span v, cuda::stream_ref stream) { using value_type = std::remove_cv_t; std::vector result(v.size()); @@ -318,7 +319,7 @@ std::vector> make_std_vector_async(device_span v, */ template std::vector make_std_vector_async(Container const& c, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) requires(std::is_convertible_v>) { return make_std_vector_async(device_span{c}, stream); @@ -336,10 +337,10 @@ std::vector make_std_vector_async(Container cons * @return The data copied to the host */ template -std::vector make_std_vector(device_span v, rmm::cuda_stream_view stream) +std::vector make_std_vector(device_span v, cuda::stream_ref stream) { auto result = make_std_vector_async(v, stream); - stream.synchronize(); + cudf::detail::sync_stream(stream); return result; } @@ -356,7 +357,7 @@ std::vector make_std_vector(device_span v, rmm::cuda_stream_view str */ template std::vector make_std_vector(Container const& c, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) requires(std::is_convertible_v>) { return make_std_vector(device_span{c}, stream); @@ -373,7 +374,7 @@ std::vector make_std_vector(Container const& c, * @return A host_vector of the given size */ template -host_vector make_host_vector(size_t size, rmm::cuda_stream_view stream) +host_vector make_host_vector(size_t size, cuda::stream_ref stream) { return host_vector(size, get_host_allocator(size, stream)); } @@ -389,7 +390,7 @@ host_vector make_host_vector(size_t size, rmm::cuda_stream_view stream) * @return A host_vector with the given capacity */ template -host_vector make_empty_host_vector(size_t capacity, rmm::cuda_stream_view stream) +host_vector make_empty_host_vector(size_t capacity, cuda::stream_ref stream) { auto result = host_vector(get_host_allocator(capacity, stream)); result.reserve(capacity); @@ -409,7 +410,7 @@ host_vector make_empty_host_vector(size_t capacity, rmm::cuda_stream_view str * @return The data copied to the host */ template -host_vector make_host_vector_async(device_span v, rmm::cuda_stream_view stream) +host_vector make_host_vector_async(device_span v, cuda::stream_ref stream) { auto result = make_host_vector(v.size(), stream); cuda_memcpy_async(result, v, stream); @@ -430,7 +431,7 @@ host_vector make_host_vector_async(device_span v, rmm::cuda_stream_v */ template host_vector make_host_vector_async(Container const& c, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) requires(std::is_convertible_v>) { return make_host_vector_async(device_span{c}, stream); @@ -449,10 +450,10 @@ host_vector make_host_vector_async(Container con * @return The data copied to the host */ template -host_vector make_host_vector(device_span v, rmm::cuda_stream_view stream) +host_vector make_host_vector(device_span v, cuda::stream_ref stream) { auto result = make_host_vector_async(v, stream); - stream.synchronize(); + cudf::detail::sync_stream(stream); return result; } @@ -469,7 +470,7 @@ host_vector make_host_vector(device_span v, rmm::cuda_stream_view st */ template host_vector make_host_vector(Container const& c, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) requires(std::is_convertible_v>) { return make_host_vector(device_span{c}, stream); @@ -484,7 +485,7 @@ host_vector make_host_vector(Container const& c, * @return A pinned host_vector with the given capacity */ template -host_vector make_empty_pinned_vector(size_t capacity, rmm::cuda_stream_view stream) +host_vector make_empty_pinned_vector(size_t capacity, cuda::stream_ref stream) { auto result = host_vector({cudf::get_pinned_memory_resource(), stream}); result.reserve(capacity); @@ -502,7 +503,7 @@ host_vector make_empty_pinned_vector(size_t capacity, rmm::cuda_stream_view s * @return A host_vector of the given size */ template -host_vector make_pinned_vector_async(size_t size, rmm::cuda_stream_view stream) +host_vector make_pinned_vector_async(size_t size, cuda::stream_ref stream) { return host_vector(size, {cudf::get_pinned_memory_resource(), stream}); } @@ -518,10 +519,10 @@ host_vector make_pinned_vector_async(size_t size, rmm::cuda_stream_view strea * @return A host_vector of the given size */ template -host_vector make_pinned_vector(size_t size, rmm::cuda_stream_view stream) +host_vector make_pinned_vector(size_t size, cuda::stream_ref stream) { auto result = make_pinned_vector_async(size, stream); - stream.synchronize(); + cudf::detail::sync_stream(stream); return result; } @@ -537,7 +538,7 @@ host_vector make_pinned_vector(size_t size, rmm::cuda_stream_view stream) * @return The data copied to pinned host memory */ template -host_vector make_pinned_vector_async(device_span v, rmm::cuda_stream_view stream) +host_vector make_pinned_vector_async(device_span v, cuda::stream_ref stream) { auto result = make_pinned_vector_async(v.size(), stream); cuda_memcpy_async(result, v, stream); @@ -557,7 +558,7 @@ host_vector make_pinned_vector_async(device_span v, rmm::cuda_stream */ template host_vector make_pinned_vector_async(Container const& c, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) requires(std::is_convertible_v>) { return make_pinned_vector_async(device_span{c}, stream); @@ -575,10 +576,10 @@ host_vector make_pinned_vector_async(Container c * @return The data copied to pinned host memory */ template -host_vector make_pinned_vector(device_span v, rmm::cuda_stream_view stream) +host_vector make_pinned_vector(device_span v, cuda::stream_ref stream) { auto result = make_pinned_vector_async(v, stream); - stream.synchronize(); + cudf::detail::sync_stream(stream); return result; } @@ -595,7 +596,7 @@ host_vector make_pinned_vector(device_span v, rmm::cuda_stream_view */ template host_vector make_pinned_vector(Container const& c, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) requires(std::is_convertible_v>) { return make_pinned_vector(device_span{c}, stream); @@ -611,7 +612,7 @@ host_vector make_pinned_vector(Container const& * @return The data copied to pinned host memory */ template -host_vector make_pinned_vector(host_span v, rmm::cuda_stream_view stream) +host_vector make_pinned_vector(host_span v, cuda::stream_ref stream) { auto result = make_pinned_vector(v.size(), stream); std::copy(v.begin(), v.end(), result.begin()); diff --git a/cpp/include/cudf/detail/valid_if.cuh b/cpp/include/cudf/detail/valid_if.cuh index 94a6e2f6a11e..486c3d080ef7 100644 --- a/cpp/include/cudf/detail/valid_if.cuh +++ b/cpp/include/cudf/detail/valid_if.cuh @@ -15,9 +15,8 @@ #include #include -#include - #include +#include namespace cudf { namespace detail { @@ -80,7 +79,7 @@ template std::pair valid_if(InputIterator begin, InputIterator end, Predicate p, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources resources) { CUDF_EXPECTS(begin <= end, "Invalid range."); @@ -97,7 +96,7 @@ std::pair valid_if(InputIterator begin, constexpr size_type block_size{256}; grid_1d grid{size, block_size}; - valid_if_kernel<<>>( + valid_if_kernel<<>>( static_cast(null_mask.data()), begin, size, p, valid_count.data()); null_count = size - valid_count.value(stream); diff --git a/cpp/include/cudf/hashing.hpp b/cpp/include/cudf/hashing.hpp index 669f4444751c..45bb40898dcc 100644 --- a/cpp/include/cudf/hashing.hpp +++ b/cpp/include/cudf/hashing.hpp @@ -53,7 +53,7 @@ namespace hashing { std::unique_ptr murmurhash3_x86_32( table_view const& input, uint32_t seed = DEFAULT_HASH_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()); /** @@ -72,7 +72,7 @@ std::unique_ptr murmurhash3_x86_32( std::unique_ptr
murmurhash3_x64_128( table_view const& input, uint64_t seed = DEFAULT_HASH_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()); /** @@ -86,7 +86,7 @@ std::unique_ptr
murmurhash3_x64_128( */ std::unique_ptr md5( 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 = cudf::get_current_device_resource_ref()); /** @@ -100,7 +100,7 @@ std::unique_ptr md5( */ std::unique_ptr sha1( 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 = cudf::get_current_device_resource_ref()); /** @@ -114,7 +114,7 @@ std::unique_ptr sha1( */ std::unique_ptr sha224( 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 = cudf::get_current_device_resource_ref()); /** @@ -128,7 +128,7 @@ std::unique_ptr sha224( */ std::unique_ptr sha256( 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 = cudf::get_current_device_resource_ref()); /** @@ -142,7 +142,7 @@ std::unique_ptr sha256( */ std::unique_ptr sha384( 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 = cudf::get_current_device_resource_ref()); /** @@ -156,7 +156,7 @@ std::unique_ptr sha384( */ std::unique_ptr sha512( 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 = cudf::get_current_device_resource_ref()); /** @@ -176,7 +176,7 @@ std::unique_ptr sha512( std::unique_ptr xxhash_32( table_view const& input, uint32_t seed = DEFAULT_HASH_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()); /** @@ -194,7 +194,7 @@ std::unique_ptr xxhash_32( std::unique_ptr xxhash_64( table_view const& input, uint64_t seed = DEFAULT_HASH_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()); /** @} */ // end of group diff --git a/cpp/include/cudf/hashing/detail/hashing.hpp b/cpp/include/cudf/hashing/detail/hashing.hpp index 567bb864628a..71b7e73e440f 100644 --- a/cpp/include/cudf/hashing/detail/hashing.hpp +++ b/cpp/include/cudf/hashing/detail/hashing.hpp @@ -8,7 +8,7 @@ #include #include -#include +#include #include #include @@ -19,46 +19,46 @@ namespace hashing::detail { std::unique_ptr murmurhash3_x86_32(table_view const& input, uint32_t seed, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref mr); std::unique_ptr
murmurhash3_x64_128(table_view const& input, uint64_t seed, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref mr); std::unique_ptr md5(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr sha1(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr sha224(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr sha256(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr sha384(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr sha512(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); std::unique_ptr xxhash_32(table_view const& input, uint64_t seed, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref mr); std::unique_ptr xxhash_64(table_view const& input, uint64_t seed, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref mr); /* SPDX-SnippetBegin diff --git a/cpp/include/cudf/interop.hpp b/cpp/include/cudf/interop.hpp index ccc000730e6d..49f904c5c871 100644 --- a/cpp/include/cudf/interop.hpp +++ b/cpp/include/cudf/interop.hpp @@ -70,7 +70,7 @@ namespace CUDF_EXPORT cudf { */ std::unique_ptr
from_dlpack( DLManagedTensor const* managed_tensor, - 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()); /** @@ -94,7 +94,7 @@ std::unique_ptr
from_dlpack( */ DLManagedTensor* to_dlpack( 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 = cudf::get_current_device_resource_ref()); /** @} */ // end of group @@ -240,7 +240,7 @@ class arrow_column { */ arrow_column(cudf::column&& input, column_metadata const& metadata, - 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()); /** @@ -257,7 +257,7 @@ class arrow_column { */ arrow_column(ArrowSchema&& schema, ArrowDeviceArray&& 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()); /** @@ -274,7 +274,7 @@ class arrow_column { */ arrow_column(ArrowSchema&& schema, ArrowArray&& 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()); /** @@ -289,7 +289,7 @@ class arrow_column { * @param mr Device memory resource used for any allocations during conversion */ arrow_column(ArrowArrayStream&& 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()); /** @@ -304,7 +304,7 @@ class arrow_column { */ void to_arrow_schema( ArrowSchema* output, - 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()) const; /** @@ -317,7 +317,7 @@ class arrow_column { */ void to_arrow(ArrowDeviceArray* output, ArrowDeviceType device_type = ARROW_DEVICE_CUDA, - 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()) const; /** @@ -358,7 +358,7 @@ class arrow_table { */ arrow_table(cudf::table&& input, std::span metadata, - 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()); /** @@ -375,7 +375,7 @@ class arrow_table { */ arrow_table(ArrowSchema&& schema, ArrowDeviceArray&& 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()); /** @@ -392,7 +392,7 @@ class arrow_table { */ arrow_table(ArrowSchema&& schema, ArrowArray&& 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()); /** @@ -407,7 +407,7 @@ class arrow_table { * @param mr Device memory resource used for any allocations during conversion */ arrow_table(ArrowArrayStream&& 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()); /** @@ -422,7 +422,7 @@ class arrow_table { */ void to_arrow_schema( ArrowSchema* output, - 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()) const; /** @@ -435,7 +435,7 @@ class arrow_table { */ void to_arrow(ArrowDeviceArray* output, ArrowDeviceType device_type = ARROW_DEVICE_CUDA, - 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()) const; /** @@ -499,7 +499,7 @@ unique_schema_t to_arrow_schema(cudf::table_view const& input, */ unique_device_array_t to_arrow_device( cudf::table&& table, - 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()); /** @@ -528,7 +528,7 @@ unique_device_array_t to_arrow_device( */ unique_device_array_t to_arrow_device( cudf::column&& col, - 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()); /** @@ -560,7 +560,7 @@ unique_device_array_t to_arrow_device( */ unique_device_array_t to_arrow_device( cudf::table_view const& table, - 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()); /** @@ -592,7 +592,7 @@ unique_device_array_t to_arrow_device( */ unique_device_array_t to_arrow_device( cudf::column_view const& col, - 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()); /** @@ -617,7 +617,7 @@ unique_device_array_t to_arrow_device( */ unique_device_array_t to_arrow_host( cudf::table_view const& table, - 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()); /** @@ -642,7 +642,7 @@ unique_device_array_t to_arrow_host( */ unique_device_array_t to_arrow_host( cudf::column_view const& col, - 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()); /** @@ -662,7 +662,7 @@ unique_device_array_t to_arrow_host( */ unique_device_array_t to_arrow_host_stringview( cudf::strings_column_view const& col, - 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()); /** @@ -685,7 +685,7 @@ unique_device_array_t to_arrow_host_stringview( std::unique_ptr from_arrow( ArrowSchema const* schema, ArrowArray 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()); /** @@ -704,7 +704,7 @@ std::unique_ptr from_arrow( std::unique_ptr from_arrow_column( ArrowSchema const* schema, ArrowArray 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()); /** @@ -730,7 +730,7 @@ std::unique_ptr from_arrow_column( std::unique_ptr
from_arrow_host( ArrowSchema const* schema, ArrowDeviceArray 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()); /** @@ -748,7 +748,7 @@ std::unique_ptr
from_arrow_host( */ std::unique_ptr
from_arrow_stream( ArrowArrayStream* 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()); /** @@ -766,7 +766,7 @@ std::unique_ptr
from_arrow_stream( */ std::unique_ptr from_arrow_stream_column( ArrowArrayStream* 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()); /** @@ -791,7 +791,7 @@ std::unique_ptr from_arrow_stream_column( std::unique_ptr from_arrow_host_column( ArrowSchema const* schema, ArrowDeviceArray 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()); /** @@ -834,7 +834,7 @@ std::unique_ptr from_arrow_host_column( unique_table_view_t from_arrow_device( ArrowSchema const* schema, ArrowDeviceArray 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()); /** @@ -872,7 +872,7 @@ unique_table_view_t from_arrow_device( unique_column_view_t from_arrow_device_column( ArrowSchema const* schema, ArrowDeviceArray 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()); /** @} */ // end of group diff --git a/cpp/include/cudf/json/json.hpp b/cpp/include/cudf/json/json.hpp index e884091808cd..83f4ca4adb57 100644 --- a/cpp/include/cudf/json/json.hpp +++ b/cpp/include/cudf/json/json.hpp @@ -159,7 +159,7 @@ std::unique_ptr get_json_object( cudf::strings_column_view const& col, cudf::string_scalar const& json_path, get_json_object_options options = get_json_object_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()); /** @} */ // end of doxygen group diff --git a/cpp/include/cudf/labeling/label_bins.hpp b/cpp/include/cudf/labeling/label_bins.hpp index 4ca8bd7cd26d..7b3d776e5f78 100644 --- a/cpp/include/cudf/labeling/label_bins.hpp +++ b/cpp/include/cudf/labeling/label_bins.hpp @@ -64,7 +64,7 @@ std::unique_ptr label_bins( inclusive left_inclusive, column_view const& right_edges, inclusive right_inclusive, - 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()); /** @} */ // end of group diff --git a/cpp/include/cudf/merge.hpp b/cpp/include/cudf/merge.hpp index 65e8d02032d7..087c209f8081 100644 --- a/cpp/include/cudf/merge.hpp +++ b/cpp/include/cudf/merge.hpp @@ -99,7 +99,7 @@ std::unique_ptr merge( std::vector const& key_cols, std::vector const& column_order, std::vector const& null_precedence = {}, - 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()); /** @} */ // end of group } // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/quantiles.hpp b/cpp/include/cudf/quantiles.hpp index 54ab981d05cb..f2d980724fc3 100644 --- a/cpp/include/cudf/quantiles.hpp +++ b/cpp/include/cudf/quantiles.hpp @@ -53,7 +53,7 @@ std::unique_ptr quantile( interpolation interp = interpolation::LINEAR, column_view const& ordered_indices = {}, bool exact = 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()); /** @@ -94,7 +94,7 @@ std::unique_ptr
quantiles( cudf::sorted is_input_sorted = sorted::NO, std::vector const& column_order = {}, std::vector const& null_precedence = {}, - 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 +123,7 @@ std::unique_ptr
quantiles( std::unique_ptr percentile_approx( tdigest::tdigest_column_view const& input, column_view const& percentiles, - 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()); /** @} */ // end of group diff --git a/cpp/include/cudf/replace.hpp b/cpp/include/cudf/replace.hpp index aa5d80cc497a..fac050208148 100644 --- a/cpp/include/cudf/replace.hpp +++ b/cpp/include/cudf/replace.hpp @@ -48,7 +48,7 @@ enum class replace_policy : bool { PRECEDING, FOLLOWING }; std::unique_ptr replace_nulls( column_view const& input, column_view const& replacement, - 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 +67,7 @@ std::unique_ptr replace_nulls( std::unique_ptr replace_nulls( column_view const& input, scalar const& replacement, - 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()); /** @@ -86,7 +86,7 @@ std::unique_ptr replace_nulls( std::unique_ptr replace_nulls( column_view const& input, replace_policy const& replace_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()); /** @@ -114,7 +114,7 @@ std::unique_ptr replace_nulls( std::unique_ptr replace_nans( column_view const& input, column_view const& replacement, - 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()); /** @@ -141,7 +141,7 @@ std::unique_ptr replace_nans( std::unique_ptr replace_nans( column_view const& input, scalar const& replacement, - 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()); /** @@ -160,7 +160,7 @@ std::unique_ptr find_and_replace_all( column_view const& input_col, column_view const& values_to_replace, column_view const& replacement_values, - 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()); /** @@ -215,7 +215,7 @@ std::unique_ptr clamp( scalar const& lo_replace, scalar const& hi, scalar const& hi_replace, - 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 +261,7 @@ std::unique_ptr clamp( column_view const& input, scalar const& lo, scalar const& hi, - 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()); /** @@ -281,7 +281,7 @@ std::unique_ptr clamp( */ std::unique_ptr normalize_nans_and_zeros( 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()); /** @@ -297,7 +297,7 @@ std::unique_ptr normalize_nans_and_zeros( * @param stream CUDA stream used for device memory operations and kernel launches */ void normalize_nans_and_zeros(mutable_column_view& in_out, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** @} */ // end of group } // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/reshape.hpp b/cpp/include/cudf/reshape.hpp index 037652096ffc..2eb89c729454 100644 --- a/cpp/include/cudf/reshape.hpp +++ b/cpp/include/cudf/reshape.hpp @@ -47,7 +47,7 @@ namespace CUDF_EXPORT cudf { */ std::unique_ptr interleave_columns( 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 = cudf::get_current_device_resource_ref()); /** @@ -72,7 +72,7 @@ std::unique_ptr interleave_columns( std::unique_ptr
tile( table_view const& input, size_type count, - 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()); /** @@ -99,7 +99,7 @@ enum class flip_endianness : bool { NO, YES }; std::unique_ptr byte_cast( column_view const& input_column, flip_endianness endian_configuration, - 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()); /** @@ -122,7 +122,7 @@ std::unique_ptr byte_cast( */ void table_to_array(table_view const& input, device_span output, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** @} */ // end of group diff --git a/cpp/include/cudf/round.hpp b/cpp/include/cudf/round.hpp index d9297e9193f5..431ba56e63fa 100644 --- a/cpp/include/cudf/round.hpp +++ b/cpp/include/cudf/round.hpp @@ -70,7 +70,7 @@ enum class rounding_method : int32_t { HALF_UP, HALF_EVEN }; column_view const& input, int32_t decimal_places = 0, rounding_method method = 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()); /** @@ -109,7 +109,7 @@ std::unique_ptr round_decimal( column_view const& input, int32_t decimal_places = 0, rounding_method method = 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()); /** @} */ // end of group diff --git a/cpp/include/cudf/stream_compaction.hpp b/cpp/include/cudf/stream_compaction.hpp index 46e923fd8f53..6242a1204f42 100644 --- a/cpp/include/cudf/stream_compaction.hpp +++ b/cpp/include/cudf/stream_compaction.hpp @@ -74,7 +74,7 @@ std::unique_ptr
drop_nulls( table_view const& input, std::vector const& keys, cudf::size_type keep_threshold, - 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()); /** @@ -107,7 +107,7 @@ std::unique_ptr
drop_nulls( std::unique_ptr
drop_nulls( table_view const& input, std::vector const& keys, - 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()); /** @@ -152,7 +152,7 @@ std::unique_ptr
drop_nans( table_view const& input, std::vector const& keys, cudf::size_type keep_threshold, - 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()); /** @@ -186,7 +186,7 @@ std::unique_ptr
drop_nans( std::unique_ptr
drop_nans( table_view const& input, std::vector const& keys, - 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()); /** @@ -214,7 +214,7 @@ std::unique_ptr
drop_nans( std::unique_ptr
apply_boolean_mask( table_view const& input, column_view const& boolean_mask, - 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()); /** @@ -242,7 +242,7 @@ std::unique_ptr
apply_boolean_mask( std::unique_ptr
apply_deletion_mask( table_view const& input, column_view const& deletion_mask, - 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()); /** @@ -291,7 +291,7 @@ std::unique_ptr
unique( std::vector const& keys, duplicate_keep_option keep, null_equality nulls_equal = 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()); /** @@ -324,7 +324,7 @@ std::unique_ptr
distinct( duplicate_keep_option keep = duplicate_keep_option::KEEP_ANY, null_equality nulls_equal = null_equality::EQUAL, nan_equality nans_equal = nan_equality::ALL_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()); /** @@ -350,7 +350,7 @@ std::unique_ptr distinct_indices( duplicate_keep_option keep = duplicate_keep_option::KEEP_ANY, null_equality nulls_equal = null_equality::EQUAL, nan_equality nans_equal = nan_equality::ALL_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()); /** @@ -386,7 +386,7 @@ std::unique_ptr
stable_distinct( duplicate_keep_option keep = duplicate_keep_option::KEEP_ANY, null_equality nulls_equal = null_equality::EQUAL, nan_equality nans_equal = nan_equality::ALL_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()); /** @@ -431,7 +431,7 @@ std::unique_ptr
stable_distinct( std::optional user_data = std::nullopt, null_aware is_null_aware = null_aware::NO, output_nullability predicate_nullability = output_nullability::PRESERVE, - 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()); /** @@ -482,7 +482,7 @@ std::vector> filter_extended( std::optional user_data = std::nullopt, null_aware is_null_aware = null_aware::NO, output_nullability predicate_nullability = output_nullability::PRESERVE, - 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()); /** @@ -509,7 +509,7 @@ std::unique_ptr
filter( table_view const& predicate_table, ast::expression const& predicate_expr, table_view const& filter_table, - 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()); /** @} */ diff --git a/cpp/include/cudf/table/equality.hpp b/cpp/include/cudf/table/equality.hpp index 2df742e405a6..36b7eb652d0b 100644 --- a/cpp/include/cudf/table/equality.hpp +++ b/cpp/include/cudf/table/equality.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -9,7 +9,7 @@ #include #include -#include +#include namespace CUDF_EXPORT cudf { /** @@ -31,7 +31,7 @@ namespace CUDF_EXPORT cudf { */ [[nodiscard]] bool tables_equal(table_view const& left, table_view const& right, - null_equality nulls_equal = null_equality::EQUAL, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + null_equality nulls_equal = null_equality::EQUAL, + cuda::stream_ref stream = cudf::get_default_stream()); } // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/table/table.hpp b/cpp/include/cudf/table/table.hpp index 891562976e3f..a6e3021a983d 100644 --- a/cpp/include/cudf/table/table.hpp +++ b/cpp/include/cudf/table/table.hpp @@ -9,7 +9,7 @@ #include #include -#include +#include #include #include @@ -47,7 +47,7 @@ class table { * @param mr Device memory resource to use for all device memory allocations */ explicit table(table const& other, - 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()); /** * @brief Moves the contents from a vector of `unique_ptr`s to columns to @@ -85,7 +85,7 @@ class table { * @param mr Device memory resource used for allocating the device memory for the new columns */ table(table_view view, - 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()); /** diff --git a/cpp/include/cudf/table/table_device_view.cuh b/cpp/include/cudf/table/table_device_view.cuh index 0b64924da53b..5cf59d7bc025 100644 --- a/cpp/include/cudf/table/table_device_view.cuh +++ b/cpp/include/cudf/table/table_device_view.cuh @@ -10,10 +10,11 @@ #include #include -#include #include #include +#include + #include #include @@ -165,7 +166,7 @@ class table_device_view : public detail::table_device_view_base> create( table_view source_view, - 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()); private: @@ -197,7 +198,7 @@ class mutable_table_device_view */ static std::unique_ptr> create(mutable_table_view source_view, - 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()); private: @@ -217,7 +218,7 @@ class mutable_table_device_view template std::pair, ColumnDeviceView*> create_column_device_views( HostTableView source_view, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); } // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/timezone.hpp b/cpp/include/cudf/timezone.hpp index 2cfe6277276c..f7b56d592add 100644 --- a/cpp/include/cudf/timezone.hpp +++ b/cpp/include/cudf/timezone.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -8,7 +8,7 @@ #include #include -#include +#include #include #include @@ -47,7 +47,7 @@ static constexpr uint32_t solar_cycle_entry_count = 2 * solar_cycle_years; std::unique_ptr
make_timezone_transition_table( std::optional tzif_dir, std::string_view timezone_name, - 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 CUDF_EXPORT cudf diff --git a/cpp/include/cudf/transpose.hpp b/cpp/include/cudf/transpose.hpp index 0649367fa4af..ef8f09b9cd0c 100644 --- a/cpp/include/cudf/transpose.hpp +++ b/cpp/include/cudf/transpose.hpp @@ -38,7 +38,7 @@ namespace CUDF_EXPORT cudf { */ std::pair, table_view> transpose( 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 = cudf::get_current_device_resource_ref()); /** @} */ // end of group diff --git a/cpp/include/cudf/utilities/prefetch.hpp b/cpp/include/cudf/utilities/prefetch.hpp index ba2b541b4fcd..5a17ac64d3a1 100644 --- a/cpp/include/cudf/utilities/prefetch.hpp +++ b/cpp/include/cudf/utilities/prefetch.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -30,7 +30,7 @@ std::atomic_bool& debug(); */ void prefetch(void const* ptr, std::size_t size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::cuda_device_id device_id = rmm::get_current_cuda_device()); /** @@ -50,7 +50,7 @@ void prefetch(void const* ptr, cudaError_t prefetch_noexcept( void const* ptr, std::size_t size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::cuda_device_id device_id = rmm::get_current_cuda_device()) noexcept; /** @@ -65,7 +65,7 @@ cudaError_t prefetch_noexcept( */ template void prefetch(rmm::device_uvector const& v, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::cuda_device_id device_id = rmm::get_current_cuda_device()) { if (v.is_empty()) { return; } diff --git a/cpp/include/cudf/utilities/roaring_bitmap.hpp b/cpp/include/cudf/utilities/roaring_bitmap.hpp index 29be57a4021d..1adf26f2329f 100644 --- a/cpp/include/cudf/utilities/roaring_bitmap.hpp +++ b/cpp/include/cudf/utilities/roaring_bitmap.hpp @@ -12,10 +12,10 @@ #include #include -#include #include #include +#include #include #include @@ -102,7 +102,7 @@ class roaring_bitmap { * * @param stream CUDA stream used for device memory operations and kernel launches */ - void materialize(rmm::cuda_stream_view stream = cudf::get_default_stream()) const; + void materialize(cuda::stream_ref stream = cudf::get_default_stream()) const; /** * @brief Returns the roaring bitmap type @@ -147,7 +147,7 @@ class roaring_bitmap { */ [[nodiscard]] std::unique_ptr contains_async( cudf::column_view const& keys, - 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()) const; /** @@ -165,7 +165,7 @@ class roaring_bitmap { */ void contains_async(cudf::column_view const& keys, cudf::mutable_column_view const& output, - rmm::cuda_stream_view stream) const; + cuda::stream_ref stream) const; private: //! Forward declaration of the opaque wrapper of cuco's roaring bitmap diff --git a/cpp/src/aggregation/aggregation.cu b/cpp/src/aggregation/aggregation.cu index d43118ef097d..3d23f447ff2f 100644 --- a/cpp/src/aggregation/aggregation.cu +++ b/cpp/src/aggregation/aggregation.cu @@ -47,7 +47,7 @@ struct identity_initializer { public: template - void operator()(mutable_column_view const& col, rmm::cuda_stream_view stream) + void operator()(mutable_column_view const& col, cuda::stream_ref stream) requires(is_supported()) { if constexpr (k == aggregation::SUM_OVERFLOW) { @@ -76,7 +76,7 @@ struct identity_initializer { } template - void operator()(mutable_column_view const& col, rmm::cuda_stream_view stream) + void operator()(mutable_column_view const& col, cuda::stream_ref stream) requires(not is_supported()) { CUDF_FAIL("Unsupported aggregation for initializing values"); @@ -86,7 +86,7 @@ struct identity_initializer { void initialize_with_identity(mutable_table_view const& table, host_span aggs, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // TODO: Initialize all the columns in a single kernel instead of invoking one // kernel per column diff --git a/cpp/src/ast/expression_parser.cpp b/cpp/src/ast/expression_parser.cpp index 854df2d611f0..3a732bf6588d 100644 --- a/cpp/src/ast/expression_parser.cpp +++ b/cpp/src/ast/expression_parser.cpp @@ -5,6 +5,7 @@ #include #include #include +#include #include #include #include @@ -54,7 +55,7 @@ expression_parser::expression_parser( cudf::table_view const& left, std::optional> right, bool has_nulls, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : _left{left}, _right{right}, @@ -75,14 +76,13 @@ expression_parser::expression_parser( expression_parser::expression_parser(expression const& expr, cudf::table_view const& table, bool has_nulls, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : expression_parser(expr, table, {}, has_nulls, stream, mr) { } -void expression_parser::move_to_device(rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) +void expression_parser::move_to_device(cuda::stream_ref stream, rmm::device_async_resource_ref mr) { std::vector sizes; std::vector data_pointers; @@ -114,7 +114,7 @@ void expression_parser::move_to_device(rmm::cuda_stream_view stream, } _device_data_buffer = rmm::device_buffer(host_data_buffer.data(), buffer_size, stream, mr); - stream.synchronize(); + cudf::detail::sync_stream(stream); // Create device pointers to components of plan auto device_data_buffer_ptr = static_cast(_device_data_buffer.data()); diff --git a/cpp/src/ast/expressions.cpp b/cpp/src/ast/expressions.cpp index af678a348909..d3d22f2b6568 100644 --- a/cpp/src/ast/expressions.cpp +++ b/cpp/src/ast/expressions.cpp @@ -109,7 +109,7 @@ auto operation::accept(detail::expression_transformer& visitor) const bool operation::may_evaluate_null(table_view const& left, table_view const& right, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { return std::any_of(operands.cbegin(), operands.cend(), @@ -135,7 +135,7 @@ std::reference_wrapper detail::predicate::accept( bool detail::predicate::may_evaluate_null(table_view const& left, table_view const& right, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { return false; } diff --git a/cpp/src/ast/jit/expressions.cpp b/cpp/src/ast/jit/expressions.cpp index feefd23afbbb..a5f766173be6 100644 --- a/cpp/src/ast/jit/expressions.cpp +++ b/cpp/src/ast/jit/expressions.cpp @@ -8,6 +8,8 @@ #include +#include + namespace cudf { namespace ast { namespace jit { @@ -31,7 +33,7 @@ std::reference_wrapper operation::accept( bool operation::may_evaluate_null(table_view const& left, table_view const& right, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { CUDF_FAIL("JIT operation is an internal expression and should not be evaluated directly", std::invalid_argument); diff --git a/cpp/src/ast/jit/expressions.hpp b/cpp/src/ast/jit/expressions.hpp index 488bf039f410..c99366705be3 100644 --- a/cpp/src/ast/jit/expressions.hpp +++ b/cpp/src/ast/jit/expressions.hpp @@ -8,6 +8,8 @@ #include #include +#include + #include #include #include @@ -86,7 +88,7 @@ struct operation : public ast::expression { [[nodiscard]] bool may_evaluate_null(table_view const& left, table_view const& right, - rmm::cuda_stream_view stream) const override; + cuda::stream_ref stream) const override; /** * @copydoc expression::accept diff --git a/cpp/src/copying/concatenate.cu b/cpp/src/copying/concatenate.cu index 4aa55cbcf6a0..42977e306feb 100644 --- a/cpp/src/copying/concatenate.cu +++ b/cpp/src/copying/concatenate.cu @@ -13,6 +13,7 @@ #include #include #include +#include #include #include #include @@ -89,7 +90,7 @@ auto create_device_views(host_span views, cuda::stream_ref st auto d_offsets = make_device_uvector_async(offsets, stream, cudf::get_current_device_resource_ref()); auto const output_size = offsets.back(); - stream.sync(); + cudf::detail::sync_stream(stream); return std::make_tuple( std::move(device_view_owners), std::move(d_views), std::move(d_offsets), output_size); diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu index 0e74026f3da6..b7451a2c23dc 100644 --- a/cpp/src/copying/contiguous_split.cu +++ b/cpp/src/copying/contiguous_split.cu @@ -13,6 +13,7 @@ #include #include #include +#include #include #include #include @@ -1366,7 +1367,7 @@ std::unique_ptr compute_splits( partition_buf_size_and_dst_buf_info->copy_to_host(); - stream.sync(); + cudf::detail::sync_stream(stream); return partition_buf_size_and_dst_buf_info; } @@ -1657,7 +1658,7 @@ std::unique_ptr chunk_iteration_state::create( d_batched_dst_buf_info[i].dst_offset -= *prior_iteration_size; }); } - stream.sync(); + cudf::detail::sync_stream(stream); return std::make_unique(std::move(d_batched_dst_buf_info), std::move(d_batch_offsets), std::move(num_batches_per_iteration), diff --git a/cpp/src/datetime/datetime_ops.cu b/cpp/src/datetime/datetime_ops.cu index d918ea424f10..4d999ec12662 100644 --- a/cpp/src/datetime/datetime_ops.cu +++ b/cpp/src/datetime/datetime_ops.cu @@ -23,10 +23,10 @@ #include #include -#include #include #include +#include #include namespace cudf { @@ -228,7 +228,7 @@ struct dispatch_round { std::unique_ptr operator()(rounding_function round_kind, rounding_frequency component, cudf::column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const requires(cudf::is_timestamp()) { @@ -273,14 +273,14 @@ struct launch_functor { launch_functor(column_view inp, mutable_column_view out) : input(inp), output(out) {} template - void operator()(rmm::cuda_stream_view stream) const + void operator()(cuda::stream_ref stream) const requires(!cudf::is_timestamp_t::value) { CUDF_FAIL("Cannot extract datetime component from non-timestamp column."); } template - void operator()(rmm::cuda_stream_view stream) const + void operator()(cuda::stream_ref stream) const requires(cudf::is_timestamp_t::value) { thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), @@ -294,7 +294,7 @@ struct launch_functor { // Create an output column by applying the functor to every element from the input column template std::unique_ptr apply_datetime_op(column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(is_timestamp(column.type()), "Column type should be timestamp"); @@ -329,7 +329,7 @@ struct add_calendrical_months_functor { template std::unique_ptr operator()(column_view timestamp_column, MonthIterator months_begin, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const requires(cudf::is_timestamp_t::value) { @@ -361,7 +361,7 @@ struct add_calendrical_months_functor { std::unique_ptr add_calendrical_months(column_view const& timestamp_column, column_view const& months_column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(is_timestamp(timestamp_column.type()), "Column type should be timestamp"); @@ -388,7 +388,7 @@ std::unique_ptr add_calendrical_months(column_view const& timestamp_colu std::unique_ptr add_calendrical_months(column_view const& timestamp_column, scalar const& months, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(is_timestamp(timestamp_column.type()), "Column type should be timestamp"); @@ -417,7 +417,7 @@ std::unique_ptr add_calendrical_months(column_view const& timestamp_colu std::unique_ptr round_general(rounding_function round_kind, rounding_frequency component, column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return cudf::type_dispatcher( @@ -425,7 +425,7 @@ std::unique_ptr round_general(rounding_function round_kind, } std::unique_ptr last_day_of_month(column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return detail::apply_datetime_op last_day_of_month(column_view const& column, } std::unique_ptr day_of_year(column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return detail::apply_datetime_op( @@ -441,21 +441,21 @@ std::unique_ptr day_of_year(column_view const& column, } std::unique_ptr is_leap_year(column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return apply_datetime_op(column, stream, mr); } std::unique_ptr days_in_month(column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return apply_datetime_op(column, stream, mr); } std::unique_ptr extract_quarter(column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return apply_datetime_op(column, stream, mr); @@ -463,7 +463,7 @@ std::unique_ptr extract_quarter(column_view const& column, std::unique_ptr extract_datetime_component(cudf::column_view const& column, datetime_component component, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { #define extract(field) \ @@ -491,7 +491,7 @@ std::unique_ptr extract_datetime_component(cudf::column_view const std::unique_ptr ceil_datetimes(column_view const& column, rounding_frequency freq, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -500,7 +500,7 @@ std::unique_ptr ceil_datetimes(column_view const& column, std::unique_ptr floor_datetimes(column_view const& column, rounding_frequency freq, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -509,7 +509,7 @@ std::unique_ptr floor_datetimes(column_view const& column, std::unique_ptr round_datetimes(column_view const& column, rounding_frequency freq, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -518,7 +518,7 @@ std::unique_ptr round_datetimes(column_view const& column, std::unique_ptr extract_datetime_component(cudf::column_view const& column, datetime_component component, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -526,7 +526,7 @@ std::unique_ptr extract_datetime_component(cudf::column_view const } std::unique_ptr last_day_of_month(column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -534,7 +534,7 @@ std::unique_ptr last_day_of_month(column_view const& column, } std::unique_ptr day_of_year(column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -543,7 +543,7 @@ std::unique_ptr day_of_year(column_view const& column, std::unique_ptr add_calendrical_months(cudf::column_view const& timestamp_column, cudf::column_view const& months_column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -552,7 +552,7 @@ std::unique_ptr add_calendrical_months(cudf::column_view const& ti std::unique_ptr add_calendrical_months(cudf::column_view const& timestamp_column, cudf::scalar const& months, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -560,7 +560,7 @@ std::unique_ptr add_calendrical_months(cudf::column_view const& ti } std::unique_ptr is_leap_year(column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -568,7 +568,7 @@ std::unique_ptr is_leap_year(column_view const& column, } std::unique_ptr days_in_month(column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -576,7 +576,7 @@ std::unique_ptr days_in_month(column_view const& column, } std::unique_ptr extract_quarter(column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/datetime/timezone.cpp b/cpp/src/datetime/timezone.cpp index 0402f476f94b..c08aa9349d7c 100644 --- a/cpp/src/datetime/timezone.cpp +++ b/cpp/src/datetime/timezone.cpp @@ -4,6 +4,7 @@ */ #include #include +#include #include #include @@ -452,7 +453,7 @@ static int64_t get_transition_time(dst_transition_s const& trans, int year) std::unique_ptr
make_timezone_transition_table(std::optional tzif_dir, std::string_view timezone_name, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -463,7 +464,7 @@ namespace detail { std::unique_ptr
make_timezone_transition_table(std::optional tzif_dir, std::string_view timezone_name, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (timezone_name == "UTC" || timezone_name.empty()) { @@ -576,7 +577,7 @@ std::unique_ptr
make_timezone_transition_table(std::optional(std::move(d_offsets), rmm::device_buffer{}, 0)); // Need to finish copies before transition_times and offsets go out of scope - stream.synchronize(); + cudf::detail::sync_stream(stream); return std::make_unique(std::move(tz_table_columns)); } diff --git a/cpp/src/groupby/streaming_groupby/impl.cu b/cpp/src/groupby/streaming_groupby/impl.cu index f7d13d97714c..7ea4569a5f06 100644 --- a/cpp/src/groupby/streaming_groupby/impl.cu +++ b/cpp/src/groupby/streaming_groupby/impl.cu @@ -89,7 +89,7 @@ streaming_groupby::impl::impl(host_span key_indices, : _max_distinct_keys{max_distinct_keys}, _null_handling{null_handling}, _mr{std::move(mr)}, - _d_agg_kinds{0, rmm::cuda_stream_default, cudf::get_current_device_resource_ref()}, + _d_agg_kinds{0, cuda::stream_ref{cudaStreamLegacy}, cudf::get_current_device_resource_ref()}, _d_agg_results{nullptr, +[](mutable_table_device_view*) {}} { CUDF_EXPECTS(max_distinct_keys > 0, "max_distinct_keys must be positive.", std::invalid_argument); diff --git a/cpp/src/hash/md5_hash.cu b/cpp/src/hash/md5_hash.cu index a00bd23f88c7..0db455ad05b3 100644 --- a/cpp/src/hash/md5_hash.cu +++ b/cpp/src/hash/md5_hash.cu @@ -17,11 +17,11 @@ #include #include -#include #include #include #include +#include #include #include @@ -274,7 +274,7 @@ inline bool md5_leaf_type_check(data_type dt) } // namespace std::unique_ptr md5(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.num_columns() == 0 || input.num_rows() == 0) { @@ -341,7 +341,7 @@ std::unique_ptr md5(table_view const& input, } // namespace detail std::unique_ptr md5(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/hash/murmurhash3_x64_128.cu b/cpp/src/hash/murmurhash3_x64_128.cu index 17af144fda06..ed1d1656c2ec 100644 --- a/cpp/src/hash/murmurhash3_x64_128.cu +++ b/cpp/src/hash/murmurhash3_x64_128.cu @@ -10,12 +10,12 @@ #include #include -#include #include #include #include #include +#include #include namespace cudf { @@ -102,7 +102,7 @@ class murmur_device_row_hasher { std::unique_ptr
murmurhash3_x64_128(table_view const& input, uint64_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto output1 = make_numeric_column( @@ -133,7 +133,7 @@ std::unique_ptr
murmurhash3_x64_128(table_view const& input, std::unique_ptr
murmurhash3_x64_128(table_view const& input, uint64_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/hash/murmurhash3_x86_32.cu b/cpp/src/hash/murmurhash3_x86_32.cu index 4e1f5ced36d5..edb6edb1877b 100644 --- a/cpp/src/hash/murmurhash3_x86_32.cu +++ b/cpp/src/hash/murmurhash3_x86_32.cu @@ -17,10 +17,10 @@ #include #include -#include #include #include +#include #include #include @@ -37,7 +37,7 @@ std::unique_ptr murmurhash3_x86_32_impl( size_type num_rows, uint32_t seed, Nullate nulls, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto output = make_numeric_column( @@ -55,7 +55,7 @@ std::unique_ptr murmurhash3_x86_32_impl( CUDF_CUDA_TRY(cub::DeviceFor::Bulk( num_rows, [output_begin, hasher] __device__(size_type i) mutable { output_begin[i] = hasher(i); }, - stream.value())); + stream.get())); return output; } @@ -64,7 +64,7 @@ std::unique_ptr murmurhash3_x86_32_impl( std::unique_ptr murmurhash3_x86_32(table_view const& input, uint32_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const preprocessed_input = cudf::detail::row::hash::preprocessed_table::create( @@ -77,7 +77,7 @@ std::unique_ptr murmurhash3_x86_32( std::shared_ptr const& input, size_type num_rows, uint32_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return murmurhash3_x86_32_impl(input, num_rows, seed, nullate::YES{}, stream, mr); @@ -87,7 +87,7 @@ std::unique_ptr murmurhash3_x86_32( std::unique_ptr murmurhash3_x86_32(table_view const& input, uint32_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/hash/murmurhash3_x86_32.cuh b/cpp/src/hash/murmurhash3_x86_32.cuh index 06d08bac8d5d..b86ee97fc958 100644 --- a/cpp/src/hash/murmurhash3_x86_32.cuh +++ b/cpp/src/hash/murmurhash3_x86_32.cuh @@ -7,9 +7,10 @@ #include -#include #include +#include + #include #include @@ -26,7 +27,7 @@ std::unique_ptr murmurhash3_x86_32( std::shared_ptr const& input, size_type num_rows, uint32_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace hashing::detail diff --git a/cpp/src/hash/sha1_hash.cu b/cpp/src/hash/sha1_hash.cu index eefff162542d..8b5b239560ac 100644 --- a/cpp/src/hash/sha1_hash.cu +++ b/cpp/src/hash/sha1_hash.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -10,9 +10,10 @@ #include #include -#include #include +#include + #include namespace cudf { @@ -50,7 +51,7 @@ struct SHA1Hash : HashBase { } // namespace std::unique_ptr sha1(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return sha_hash(input, stream, mr); @@ -59,7 +60,7 @@ std::unique_ptr sha1(table_view const& input, } // namespace detail std::unique_ptr sha1(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/hash/sha224_hash.cu b/cpp/src/hash/sha224_hash.cu index c5f760463756..b4cc139a3cad 100644 --- a/cpp/src/hash/sha224_hash.cu +++ b/cpp/src/hash/sha224_hash.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -10,9 +10,10 @@ #include #include -#include #include +#include + #include namespace cudf { @@ -51,7 +52,7 @@ struct SHA224Hash : HashBase { } // namespace std::unique_ptr sha224(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return sha_hash(input, stream, mr); @@ -60,7 +61,7 @@ std::unique_ptr sha224(table_view const& input, } // namespace detail std::unique_ptr sha224(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/hash/sha256_hash.cu b/cpp/src/hash/sha256_hash.cu index 747feb55263c..2bb6a269e0d7 100644 --- a/cpp/src/hash/sha256_hash.cu +++ b/cpp/src/hash/sha256_hash.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -10,9 +10,10 @@ #include #include -#include #include +#include + #include namespace cudf { @@ -51,7 +52,7 @@ struct SHA256Hash : HashBase { } // namespace std::unique_ptr sha256(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return sha_hash(input, stream, mr); @@ -60,7 +61,7 @@ std::unique_ptr sha256(table_view const& input, } // namespace detail std::unique_ptr sha256(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/hash/sha384_hash.cu b/cpp/src/hash/sha384_hash.cu index f25519311087..266643c05210 100644 --- a/cpp/src/hash/sha384_hash.cu +++ b/cpp/src/hash/sha384_hash.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -10,9 +10,10 @@ #include #include -#include #include +#include + #include namespace cudf { @@ -58,7 +59,7 @@ struct SHA384Hash : HashBase { } // namespace std::unique_ptr sha384(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return sha_hash(input, stream, mr); @@ -67,7 +68,7 @@ std::unique_ptr sha384(table_view const& input, } // namespace detail std::unique_ptr sha384(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/hash/sha512_hash.cu b/cpp/src/hash/sha512_hash.cu index 03f4a9a02db6..4d2c01bb41c8 100644 --- a/cpp/src/hash/sha512_hash.cu +++ b/cpp/src/hash/sha512_hash.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -10,9 +10,10 @@ #include #include -#include #include +#include + #include namespace cudf { @@ -58,7 +59,7 @@ struct SHA512Hash : HashBase { } // namespace std::unique_ptr sha512(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return sha_hash(input, stream, mr); @@ -67,7 +68,7 @@ std::unique_ptr sha512(table_view const& input, } // namespace detail std::unique_ptr sha512(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/hash/sha_hash.cuh b/cpp/src/hash/sha_hash.cuh index 0344acad5331..e2189f9fbf07 100644 --- a/cpp/src/hash/sha_hash.cuh +++ b/cpp/src/hash/sha_hash.cuh @@ -19,11 +19,11 @@ #include #include -#include #include #include #include +#include #include #include #include @@ -495,7 +495,7 @@ bool inline sha_leaf_type_check(data_type dt) */ template std::unique_ptr sha_hash(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.num_rows() == 0) { return cudf::make_empty_column(cudf::type_id::STRING); } diff --git a/cpp/src/hash/xxhash_32.cu b/cpp/src/hash/xxhash_32.cu index 0efd6f02ce12..ca31b0d26961 100644 --- a/cpp/src/hash/xxhash_32.cu +++ b/cpp/src/hash/xxhash_32.cu @@ -9,9 +9,8 @@ #include #include -#include - #include +#include namespace cudf { namespace hashing { @@ -19,7 +18,7 @@ namespace detail { std::unique_ptr xxhash_32(table_view const& input, uint32_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto output = make_numeric_column(data_type(type_to_id()), @@ -42,7 +41,7 @@ std::unique_ptr xxhash_32(table_view const& input, CUDF_CUDA_TRY(cub::DeviceFor::Bulk( input.num_rows(), [output_begin, hasher] __device__(size_type i) mutable { output_begin[i] = hasher(i); }, - stream.value())); + stream.get())); return output; } @@ -51,7 +50,7 @@ std::unique_ptr xxhash_32(table_view const& input, std::unique_ptr xxhash_32(table_view const& input, uint32_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/hash/xxhash_64.cu b/cpp/src/hash/xxhash_64.cu index 69c97724bbd6..35a66008077f 100644 --- a/cpp/src/hash/xxhash_64.cu +++ b/cpp/src/hash/xxhash_64.cu @@ -9,9 +9,8 @@ #include #include -#include - #include +#include namespace cudf { namespace hashing { @@ -21,7 +20,7 @@ using hash_value_type = uint64_t; std::unique_ptr xxhash_64(table_view const& input, uint64_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto output = make_numeric_column(data_type(type_to_id()), @@ -44,7 +43,7 @@ std::unique_ptr xxhash_64(table_view const& input, CUDF_CUDA_TRY(cub::DeviceFor::Bulk( input.num_rows(), [output_begin, hasher] __device__(size_type i) mutable { output_begin[i] = hasher(i); }, - stream.value())); + stream.get())); return output; } @@ -53,7 +52,7 @@ std::unique_ptr xxhash_64(table_view const& input, std::unique_ptr xxhash_64(table_view const& input, uint64_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/interop/arrow_data_structures.cpp b/cpp/src/interop/arrow_data_structures.cpp index 63ffead8641f..12522201fbab 100644 --- a/cpp/src/interop/arrow_data_structures.cpp +++ b/cpp/src/interop/arrow_data_structures.cpp @@ -9,9 +9,10 @@ #include #include -#include #include +#include + #include #include #include @@ -70,7 +71,7 @@ struct arrow_array_container { template arrow_array_container(ArrowSchema&& schema_, T input_, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto output = cudf::to_arrow_device(std::move(input_), stream, mr); @@ -80,7 +81,7 @@ struct arrow_array_container { arrow_array_container(ArrowSchema&& schema_, ArrowDeviceArray&& input_, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(!contains_fixed_size_list(schema_), @@ -201,7 +202,7 @@ void arrow_obj_to_arrow(T& obj, std::shared_ptr container, ArrowDeviceArray* output, ArrowDeviceType device_type, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { switch (device_type) { @@ -234,7 +235,7 @@ void arrow_obj_to_arrow(T& obj, arrow_column::arrow_column(cudf::column&& input, column_metadata const& metadata, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : container{[&] { auto table_meta = std::vector{metadata}; @@ -251,7 +252,7 @@ arrow_column::arrow_column(cudf::column&& input, arrow_column::arrow_column(ArrowSchema&& schema, ArrowDeviceArray&& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { switch (input.device_type) { @@ -275,7 +276,7 @@ arrow_column::arrow_column(ArrowSchema&& schema, arrow_column::arrow_column(ArrowSchema&& schema, ArrowArray&& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { ArrowDeviceArray arr{.array = {}, .device_id = -1, .device_type = ARROW_DEVICE_CPU}; @@ -287,7 +288,7 @@ arrow_column::arrow_column(ArrowSchema&& schema, } arrow_column::arrow_column(ArrowArrayStream&& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto col = from_arrow_stream_column(&input, stream, mr); @@ -298,7 +299,7 @@ arrow_column::arrow_column(ArrowArrayStream&& input, } void arrow_column::to_arrow_schema(ArrowSchema* output, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { NANOARROW_THROW_NOT_OK(ArrowSchemaDeepCopy(&container->schema, output)); @@ -306,7 +307,7 @@ void arrow_column::to_arrow_schema(ArrowSchema* output, void arrow_column::to_arrow(ArrowDeviceArray* output, ArrowDeviceType device_type, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { arrow_obj_to_arrow(*this, container, output, device_type, stream, mr); @@ -316,7 +317,7 @@ column_view arrow_column::view() const { return cached_view; } arrow_table::arrow_table(cudf::table&& input, std::span metadata, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : container{[&]() { auto schema = cudf::to_arrow_schema(input.view(), metadata); @@ -331,7 +332,7 @@ arrow_table::arrow_table(cudf::table&& input, arrow_table::arrow_table(ArrowSchema&& schema, ArrowDeviceArray&& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { switch (input.device_type) { @@ -357,7 +358,7 @@ arrow_table::arrow_table(ArrowSchema&& schema, arrow_table::arrow_table(ArrowSchema&& schema, ArrowArray&& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { ArrowDeviceArray arr{.array = {}, .device_id = -1, .device_type = ARROW_DEVICE_CPU}; @@ -369,7 +370,7 @@ arrow_table::arrow_table(ArrowSchema&& schema, } arrow_table::arrow_table(ArrowArrayStream&& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto tbl = from_arrow_stream(&input, stream, mr); @@ -380,7 +381,7 @@ arrow_table::arrow_table(ArrowArrayStream&& input, } void arrow_table::to_arrow_schema(ArrowSchema* output, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { NANOARROW_THROW_NOT_OK(ArrowSchemaDeepCopy(&container->schema, output)); @@ -388,7 +389,7 @@ void arrow_table::to_arrow_schema(ArrowSchema* output, void arrow_table::to_arrow(ArrowDeviceArray* output, ArrowDeviceType device_type, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { arrow_obj_to_arrow(*this, container, output, device_type, stream, mr); diff --git a/cpp/src/interop/dlpack.cpp b/cpp/src/interop/dlpack.cpp index e29036890f5b..6e1d2bfb87f2 100644 --- a/cpp/src/interop/dlpack.cpp +++ b/cpp/src/interop/dlpack.cpp @@ -1,10 +1,11 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include #include #include +#include #include #include #include @@ -12,7 +13,7 @@ #include #include -#include +#include #include @@ -123,7 +124,7 @@ struct dltensor_context { namespace detail { std::unique_ptr
from_dlpack(DLManagedTensor const* managed_tensor, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(nullptr != managed_tensor, "managed_tensor is null"); @@ -206,7 +207,7 @@ std::unique_ptr
from_dlpack(DLManagedTensor const* managed_tensor, } DLManagedTensor* to_dlpack(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const num_rows = input.num_rows(); @@ -274,7 +275,7 @@ DLManagedTensor* to_dlpack(table_view const& input, // synchronize the stream because after the return the data may be accessed from the host before // the above async copies have completed (especially if pinned host memory is used). - stream.synchronize(); + cudf::detail::sync_stream(stream); return managed_tensor.release(); } @@ -282,7 +283,7 @@ DLManagedTensor* to_dlpack(table_view const& input, } // namespace detail std::unique_ptr
from_dlpack(DLManagedTensor const* managed_tensor, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -290,7 +291,7 @@ std::unique_ptr
from_dlpack(DLManagedTensor const* managed_tensor, } DLManagedTensor* to_dlpack(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/interop/from_arrow_device.cu b/cpp/src/interop/from_arrow_device.cu index bfa28026bd06..3a79b001bad7 100644 --- a/cpp/src/interop/from_arrow_device.cu +++ b/cpp/src/interop/from_arrow_device.cu @@ -14,6 +14,7 @@ #include #include #include +#include #include #include #include @@ -24,10 +25,10 @@ #include #include -#include #include #include +#include #include #include @@ -51,7 +52,7 @@ struct dispatch_from_arrow_device { ArrowArray const*, data_type, bool, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref) { CUDF_FAIL("Unsupported type in from_arrow_device", cudf::data_type_error); @@ -62,7 +63,7 @@ struct dispatch_from_arrow_device { ArrowArray const* input, data_type type, bool skip_mask, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref mr) { size_type const num_rows = input->length; @@ -84,7 +85,7 @@ dispatch_tuple_t get_column(ArrowSchemaView* schema, ArrowArray const* input, data_type type, bool skip_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); template <> @@ -92,7 +93,7 @@ dispatch_tuple_t dispatch_from_arrow_device::operator()(ArrowSchemaView* s ArrowArray const* input, data_type type, bool skip_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input->length == 0) { @@ -155,7 +156,7 @@ dispatch_tuple_t dispatch_from_arrow_device::operator()( ArrowArray const* input, data_type type, bool skip_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto null_mask = skip_mask @@ -197,7 +198,7 @@ dispatch_tuple_t dispatch_from_arrow_device::operator()( // gather strings into output column auto out_col = cudf::make_strings_column(d_indices, stream, mr); owned.emplace_back(std::move(out_col)); - stream.synchronize(); + cudf::detail::sync_stream(stream); return std::make_tuple(owned.front()->view(), std::move(owned)); } @@ -222,7 +223,7 @@ dispatch_tuple_t dispatch_from_arrow_device::operator()( ArrowArray const* input, data_type type, bool skip_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { ArrowSchemaView keys_schema_view; @@ -271,7 +272,7 @@ dispatch_tuple_t dispatch_from_arrow_device::operator()( ArrowArray const* input, data_type type, bool skip_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { std::vector children; @@ -316,7 +317,7 @@ dispatch_tuple_t dispatch_from_arrow_device::operator()( ArrowArray const* input, data_type type, bool skip_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(schema->type != NANOARROW_TYPE_LARGE_LIST, @@ -394,7 +395,7 @@ dispatch_tuple_t get_column(ArrowSchemaView* schema, ArrowArray const* input, data_type type, bool skip_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS( @@ -417,7 +418,7 @@ dispatch_tuple_t get_column(ArrowSchemaView* schema, unique_table_view_t from_arrow_device(ArrowSchema const* schema, ArrowDeviceArray const* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(schema != nullptr && input != nullptr, @@ -436,7 +437,7 @@ unique_table_view_t from_arrow_device(ArrowSchema const* schema, if (input->sync_event != nullptr) { CUDF_CUDA_TRY( - cudaStreamWaitEvent(stream.value(), *reinterpret_cast(input->sync_event))); + cudaStreamWaitEvent(stream.get(), *reinterpret_cast(input->sync_event))); } std::vector columns; @@ -484,7 +485,7 @@ unique_table_view_t from_arrow_device(ArrowSchema const* schema, unique_column_view_t from_arrow_device_column(ArrowSchema const* schema, ArrowDeviceArray const* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(schema != nullptr && input != nullptr, @@ -503,7 +504,7 @@ unique_column_view_t from_arrow_device_column(ArrowSchema const* schema, if (input->sync_event != nullptr) { CUDF_CUDA_TRY( - cudaStreamWaitEvent(stream.value(), *reinterpret_cast(input->sync_event))); + cudaStreamWaitEvent(stream.get(), *reinterpret_cast(input->sync_event))); } auto type = arrow_to_cudf_type(&view); @@ -516,7 +517,7 @@ unique_column_view_t from_arrow_device_column(ArrowSchema const* schema, unique_table_view_t from_arrow_device(ArrowSchema const* schema, ArrowDeviceArray const* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -526,7 +527,7 @@ unique_table_view_t from_arrow_device(ArrowSchema const* schema, unique_column_view_t from_arrow_device_column(ArrowSchema const* schema, ArrowDeviceArray const* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/interop/from_arrow_host.cu b/cpp/src/interop/from_arrow_host.cu index f777b7b6aba5..22c712f2b151 100644 --- a/cpp/src/interop/from_arrow_host.cu +++ b/cpp/src/interop/from_arrow_host.cu @@ -12,6 +12,7 @@ #include #include #include +#include #include #include #include @@ -26,10 +27,10 @@ #include #include -#include #include #include +#include #include #include @@ -93,7 +94,7 @@ CUDF_KERNEL void copy_shifted_bitmask(bitmask_type* __restrict__ destination, // copies the bitmask to device and automatically applies the offset std::pair, size_type> get_mask_buffer( - ArrowArray const* input, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + ArrowArray const* input, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input->length == 0) { return {std::make_unique(0, stream, mr), 0}; } @@ -117,9 +118,9 @@ std::pair, size_type> get_mask_buffer( if (mask_words > 0 && bit_index > 0) { auto dest_mask = rmm::device_uvector(padded_words, stream, mr); cudf::detail::grid_1d config(mask_words, 256); - copy_shifted_bitmask<<>>( + copy_shifted_bitmask<<>>( dest_mask.data(), mask.data(), bit_index, bit_index + num_rows, mask_words); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); mask = std::move(dest_mask); } @@ -133,11 +134,11 @@ std::unique_ptr get_column_copy(ArrowSchemaView const* schema, ArrowArray const* input, data_type type, bool skip_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); struct dispatch_copy_from_arrow_host { - rmm::cuda_stream_view stream; + cuda::stream_ref stream; rmm::device_async_resource_ref mr; template () && !is_fixed_point())> @@ -197,9 +198,9 @@ std::unique_ptr dispatch_copy_from_arrow_host::operator()(ArrowSch if (data_words > 0 && bit_index > 0) { auto dest_data = rmm::device_uvector(data_words, stream, mr); cudf::detail::grid_1d config(data_words, 256); - copy_shifted_bitmask<<>>( + copy_shifted_bitmask<<>>( dest_data.data(), data.data(), bit_index, bit_index + num_rows, data_words); - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); data = std::move(dest_data); } @@ -374,7 +375,7 @@ std::unique_ptr get_column_copy(ArrowSchemaView const* schema, ArrowArray const* input, data_type type, bool skip_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS( @@ -407,7 +408,7 @@ template std::tuple, int64_t, int64_t> copy_offsets_column( ArrowSchemaView const* schema, ArrowArray const* offsets, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto offsets_buffer = @@ -457,7 +458,7 @@ std::unique_ptr make_fixed_size_list_offsets(size_type num_offsets, std::tuple, int64_t, int64_t> get_offsets_column( ArrowSchemaView const* schema, ArrowArray const* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { void const* offsets_buffer = input->buffers[fixed_width_data_buffer_idx]; @@ -501,13 +502,13 @@ std::tuple, int64_t, int64_t> get_offsets_column( offsets_array.offset = 0; // already accounted for by the above transform auto result = dispatch_copy_from_arrow_host{stream, mr}.template operator()( schema, &offsets_array, data_type(type_id::INT32), true); - stream.synchronize(); + cudf::detail::sync_stream(stream); return std::tuple{std::move(result), offset, length}; } std::unique_ptr
from_arrow_host(ArrowSchema const* schema, ArrowDeviceArray const* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(schema != nullptr && input != nullptr, @@ -546,13 +547,13 @@ std::unique_ptr
from_arrow_host(ArrowSchema const* schema, std::overflow_error); return std::make_unique
(std::move(columns), static_cast(input->array.length)); } - stream.synchronize(); + cudf::detail::sync_stream(stream); return std::make_unique
(std::move(columns)); } std::unique_ptr from_arrow_host_column(ArrowSchema const* schema, ArrowDeviceArray const* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(schema != nullptr && input != nullptr, @@ -567,7 +568,7 @@ std::unique_ptr from_arrow_host_column(ArrowSchema const* schema, auto type = arrow_to_cudf_type(&view); auto result = get_column_copy(&view, &input->array, type, false, stream, mr); - stream.synchronize(); + cudf::detail::sync_stream(stream); return result; } @@ -575,7 +576,7 @@ std::unique_ptr get_column_from_host_copy(ArrowSchemaView const* schema, ArrowArray const* input, data_type type, bool skip_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return get_column_copy(schema, input, type, skip_mask, stream, mr); @@ -585,7 +586,7 @@ std::unique_ptr get_column_from_host_copy(ArrowSchemaView const* schema, std::unique_ptr
from_arrow_host(ArrowSchema const* schema, ArrowDeviceArray const* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -595,7 +596,7 @@ std::unique_ptr
from_arrow_host(ArrowSchema const* schema, std::unique_ptr from_arrow_host_column(ArrowSchema const* schema, ArrowDeviceArray const* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -605,7 +606,7 @@ std::unique_ptr from_arrow_host_column(ArrowSchema const* schema, std::unique_ptr
from_arrow(ArrowSchema const* schema, ArrowArray const* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -620,7 +621,7 @@ std::unique_ptr
from_arrow(ArrowSchema const* schema, std::unique_ptr from_arrow_column(ArrowSchema const* schema, ArrowArray const* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/interop/from_arrow_host.hpp b/cpp/src/interop/from_arrow_host.hpp index 5d635ddcd179..a4991dc3c1b6 100644 --- a/cpp/src/interop/from_arrow_host.hpp +++ b/cpp/src/interop/from_arrow_host.hpp @@ -6,10 +6,11 @@ #include -#include #include #include +#include + #include #include @@ -31,7 +32,7 @@ std::unique_ptr string_column_from_arrow_host(ArrowSchemaView const* sch ArrowArray const* input, std::unique_ptr&& mask, size_type null_count, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -47,7 +48,7 @@ std::unique_ptr string_column_from_arrow_host(ArrowSchemaView const* sch std::tuple, int64_t, int64_t> get_offsets_column( ArrowSchemaView const* schema, ArrowArray const* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** diff --git a/cpp/src/interop/from_arrow_host_strings.cu b/cpp/src/interop/from_arrow_host_strings.cu index b15b8960eec1..fbc6869c84d2 100644 --- a/cpp/src/interop/from_arrow_host_strings.cu +++ b/cpp/src/interop/from_arrow_host_strings.cu @@ -10,6 +10,7 @@ #include #include #include +#include #include #include #include @@ -22,11 +23,11 @@ #include #include -#include #include #include #include +#include #include #include @@ -46,7 +47,7 @@ std::unique_ptr from_arrow_string(ArrowSchemaView const* schema, ArrowArray const* input, std::unique_ptr&& mask, size_type null_count, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto [offsets_column, offset, char_data_length] = get_offsets_column(schema, input, stream, mr); @@ -67,7 +68,7 @@ constexpr int stringview_vector_idx = 1; std::unique_ptr from_arrow_stringview(ArrowSchemaView const* schema, ArrowArray const* input, std::unique_ptr&& mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { ArrowArrayView view; @@ -113,7 +114,7 @@ std::unique_ptr from_arrow_stringview(ArrowSchemaView const* schema, return {data, size}; }); - stream.synchronize(); + cudf::detail::sync_stream(stream); return cudf::make_strings_column(d_indices, stream, mr); } @@ -123,7 +124,7 @@ std::unique_ptr string_column_from_arrow_host(ArrowSchemaView const* sch ArrowArray const* input, std::unique_ptr&& mask, size_type null_count, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return schema->type == NANOARROW_TYPE_STRING_VIEW diff --git a/cpp/src/interop/from_arrow_stream.cu b/cpp/src/interop/from_arrow_stream.cu index c6825d503dc5..1c15c68aa1ab 100644 --- a/cpp/src/interop/from_arrow_stream.cu +++ b/cpp/src/interop/from_arrow_stream.cu @@ -8,12 +8,14 @@ #include #include #include +#include #include #include -#include #include +#include + #include #include @@ -28,7 +30,7 @@ namespace detail { namespace { std::unique_ptr make_empty_column_from_schema(ArrowSchema const* schema, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { ArrowSchemaView schema_view; @@ -66,7 +68,7 @@ std::unique_ptr make_empty_column_from_schema(ArrowSchema const* schema, } // namespace std::unique_ptr
from_arrow_stream(ArrowArrayStream* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input != nullptr, "input ArrowArrayStream must not be NULL", std::invalid_argument); @@ -114,7 +116,7 @@ std::unique_ptr
from_arrow_stream(ArrowArrayStream* input, // Ensure all host-to-device copies enqueued above have completed before `sources` releases the // host-side Arrow buffers. - stream.synchronize(); + cudf::detail::sync_stream(stream); if (chunks.size() == 1) { return std::move(chunks[0]); } auto chunk_views = std::vector{}; @@ -127,7 +129,7 @@ std::unique_ptr
from_arrow_stream(ArrowArrayStream* input, } std::unique_ptr from_arrow_stream_column(ArrowArrayStream* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input != nullptr, "input ArrowArrayStream must not be NULL", std::invalid_argument); @@ -163,7 +165,7 @@ std::unique_ptr from_arrow_stream_column(ArrowArrayStream* input, // Ensure all host-to-device copies enqueued above have completed before `sources` releases the // host-side Arrow buffers. - stream.synchronize(); + cudf::detail::sync_stream(stream); if (chunks.size() == 1) { return std::move(chunks[0]); } auto chunk_views = std::vector{}; @@ -178,7 +180,7 @@ std::unique_ptr from_arrow_stream_column(ArrowArrayStream* input, } // namespace detail std::unique_ptr
from_arrow_stream(ArrowArrayStream* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -186,7 +188,7 @@ std::unique_ptr
from_arrow_stream(ArrowArrayStream* input, } std::unique_ptr from_arrow_stream_column(ArrowArrayStream* input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/interop/to_arrow_device.cu b/cpp/src/interop/to_arrow_device.cu index 2dc8bfe48c2c..27d84fafeae8 100644 --- a/cpp/src/interop/to_arrow_device.cu +++ b/cpp/src/interop/to_arrow_device.cu @@ -24,9 +24,9 @@ #include #include -#include #include +#include #include #include @@ -97,7 +97,7 @@ int set_contents(column::contents& contents, ArrowArray* out) struct dispatch_to_arrow_device { template () and not is_fixed_point())> - int operator()(cudf::column&&, rmm::cuda_stream_view, rmm::device_async_resource_ref, ArrowArray*) + int operator()(cudf::column&&, cuda::stream_ref, rmm::device_async_resource_ref, ArrowArray*) { CUDF_FAIL("Unsupported type for to_arrow_device", cudf::data_type_error); } @@ -105,7 +105,7 @@ struct dispatch_to_arrow_device { // cover rep layout compatible and decimal types template () or is_fixed_point())> int operator()(cudf::column&& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr, ArrowArray* out) { @@ -132,7 +132,7 @@ int handle_empty_type_column(ArrowArray* array, cudf::column& column) template <> int dispatch_to_arrow_device::operator()(cudf::column&& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr, ArrowArray* out) { @@ -151,7 +151,7 @@ int dispatch_to_arrow_device::operator()(cudf::column&& column, template <> int dispatch_to_arrow_device::operator()(cudf::column&& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr, ArrowArray* out) { @@ -197,19 +197,19 @@ int dispatch_to_arrow_device::operator()(cudf::column&& colum template <> int dispatch_to_arrow_device::operator()(cudf::column&& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr, ArrowArray* out); template <> int dispatch_to_arrow_device::operator()(cudf::column&& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr, ArrowArray* out); template <> int dispatch_to_arrow_device::operator()(cudf::column&& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr, ArrowArray* out) { @@ -237,7 +237,7 @@ int dispatch_to_arrow_device::operator()(cudf::column&& colum template <> int dispatch_to_arrow_device::operator()(cudf::column&& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr, ArrowArray* out) { @@ -266,7 +266,7 @@ int dispatch_to_arrow_device::operator()(cudf::column&& column, template <> int dispatch_to_arrow_device::operator()(cudf::column&& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr, ArrowArray* out) { @@ -302,7 +302,7 @@ int dispatch_to_arrow_device::operator()(cudf::column&& colu struct dispatch_to_arrow_device_view { cudf::column_view column; - rmm::cuda_stream_view stream; + cuda::stream_ref stream; rmm::device_async_resource_ref mr; template release = nullptr; } -unique_device_array_t create_device_array(nanoarrow::UniqueArray&& out, - rmm::cuda_stream_view stream) +unique_device_array_t create_device_array(nanoarrow::UniqueArray&& out, cuda::stream_ref stream) { NANOARROW_THROW_NOT_OK( ArrowArrayFinishBuilding(out.get(), NANOARROW_VALIDATION_LEVEL_MINIMAL, nullptr)); auto private_data = std::make_unique(); CUDF_CUDA_TRY(cudaEventCreate(&private_data->sync_event)); - CUDF_CUDA_TRY(cudaEventRecord(private_data->sync_event, stream.value())); + CUDF_CUDA_TRY(cudaEventRecord(private_data->sync_event, stream.get())); ArrowArrayMove(out.get(), &private_data->parent); unique_device_array_t result(new ArrowDeviceArray, [](ArrowDeviceArray* arr) { @@ -525,7 +524,7 @@ unique_device_array_t create_device_array(nanoarrow::UniqueArray&& out, } // namespace unique_device_array_t to_arrow_device(cudf::table&& table, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { nanoarrow::UniqueArray tmp; @@ -551,7 +550,7 @@ unique_device_array_t to_arrow_device(cudf::table&& table, } unique_device_array_t to_arrow_device(cudf::column&& col, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { nanoarrow::UniqueArray tmp; @@ -567,7 +566,7 @@ unique_device_array_t to_arrow_device(cudf::column&& col, } unique_device_array_t to_arrow_device(cudf::table_view const& table, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { nanoarrow::UniqueArray tmp; @@ -588,7 +587,7 @@ unique_device_array_t to_arrow_device(cudf::table_view const& table, } unique_device_array_t to_arrow_device(cudf::column_view const& col, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { nanoarrow::UniqueArray tmp; @@ -602,7 +601,7 @@ unique_device_array_t to_arrow_device(cudf::column_view const& col, } // namespace detail unique_device_array_t to_arrow_device(cudf::table&& table, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -610,7 +609,7 @@ unique_device_array_t to_arrow_device(cudf::table&& table, } unique_device_array_t to_arrow_device(cudf::column&& col, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -618,7 +617,7 @@ unique_device_array_t to_arrow_device(cudf::column&& col, } unique_device_array_t to_arrow_device(cudf::table_view const& table, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -626,7 +625,7 @@ unique_device_array_t to_arrow_device(cudf::table_view const& table, } unique_device_array_t to_arrow_device(cudf::column_view const& col, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/interop/to_arrow_host.cu b/cpp/src/interop/to_arrow_host.cu index 77685bb281f1..2b20f217922a 100644 --- a/cpp/src/interop/to_arrow_host.cu +++ b/cpp/src/interop/to_arrow_host.cu @@ -12,6 +12,7 @@ #include #include #include +#include #include #include #include @@ -27,13 +28,13 @@ #include #include -#include #include #include #include #include #include +#include #include #include @@ -75,7 +76,7 @@ void enable_hugepage(ArrowBuffer* buffer) struct dispatch_to_arrow_host { cudf::column_view column; - rmm::cuda_stream_view stream; + cuda::stream_ref stream; rmm::device_async_resource_ref mr; int populate_validity_bitmap(ArrowBitmap* bitmap) const @@ -90,7 +91,8 @@ struct dispatch_to_arrow_host { : column.null_mask(), bitmap->buffer.size_bytes, stream)); - stream.synchronize(); // ensures the bitmap is not destroyed before the copy is completed + cudf::detail::sync_stream( + stream); // ensures the bitmap is not destroyed before the copy is completed return NANOARROW_OK; } @@ -131,7 +133,7 @@ struct dispatch_to_arrow_host { }; int get_column(cudf::column_view column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr, ArrowArray* out); @@ -299,7 +301,7 @@ int dispatch_to_arrow_host::operator()(ArrowArray* out) const } int get_column(cudf::column_view column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr, ArrowArray* out) { @@ -331,7 +333,7 @@ unique_device_array_t create_device_array(nanoarrow::UniqueArray&& out) } // namespace unique_device_array_t to_arrow_host(cudf::table_view const& table, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { nanoarrow::UniqueArray tmp; @@ -351,13 +353,13 @@ unique_device_array_t to_arrow_host(cudf::table_view const& table, // wait for all the stream operations to complete before we return. // this ensures that the host memory that we're returning will be populated // before we return from this function. - stream.synchronize(); + cudf::detail::sync_stream(stream); return create_device_array(std::move(tmp)); } unique_device_array_t to_arrow_host(cudf::column_view const& col, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { nanoarrow::UniqueArray tmp; @@ -368,7 +370,7 @@ unique_device_array_t to_arrow_host(cudf::column_view const& col, // wait for all the stream operations to complete before we return. // this ensures that the host memory that we're returning will be populated // before we return from this function. - stream.synchronize(); + cudf::detail::sync_stream(stream); return create_device_array(std::move(tmp)); } @@ -417,7 +419,7 @@ struct strings_to_binary_view { }; unique_device_array_t to_arrow_host_stringview(cudf::strings_column_view const& col, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { nanoarrow::UniqueArray out; @@ -462,7 +464,7 @@ unique_device_array_t to_arrow_host_stringview(cudf::strings_column_view const& })); auto longer_strings = cudf::strings::detail::make_strings_column( indices, indices + col.size(), stream, cudf::get_current_device_resource_ref()); - stream.synchronize(); + cudf::detail::sync_stream(stream); auto const sv = cudf::strings_column_view(longer_strings->view()); return std::pair{std::move(longer_strings), sv}; }(); @@ -534,13 +536,13 @@ unique_device_array_t to_arrow_host_stringview(cudf::strings_column_view const& out->null_count = col.null_count(); out->offset = 0; - stream.synchronize(); + cudf::detail::sync_stream(stream); return create_device_array(std::move(out)); } } // namespace detail unique_device_array_t to_arrow_host(cudf::column_view const& col, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -548,7 +550,7 @@ unique_device_array_t to_arrow_host(cudf::column_view const& col, } unique_device_array_t to_arrow_host(cudf::table_view const& table, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -556,7 +558,7 @@ unique_device_array_t to_arrow_host(cudf::table_view const& table, } unique_device_array_t to_arrow_host_stringview(cudf::strings_column_view const& col, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/jit/cache.hpp b/cpp/src/jit/cache.hpp index 0df58eae95e8..f96d58c587af 100644 --- a/cpp/src/jit/cache.hpp +++ b/cpp/src/jit/cache.hpp @@ -6,7 +6,7 @@ #pragma once #include -#include +#include #include @@ -53,17 +53,17 @@ struct [[nodiscard]] kernel { void launch(rtcx::cuda_dim3 grid_dim, rtcx::cuda_dim3 block_dim, uint32_t shared_mem_bytes, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, void** kernel_params) const { - return _kernel.launch(grid_dim, block_dim, shared_mem_bytes, stream.value(), kernel_params); + return _kernel.launch(grid_dim, block_dim, shared_mem_bytes, stream.get(), kernel_params); } template void launch_with(rtcx::cuda_dim3 grid_dim, rtcx::cuda_dim3 block_dim, uint32_t shared_mem_bytes, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, Args&&... args) requires(sizeof...(Args) > 0) { diff --git a/cpp/src/jit/helpers.hpp b/cpp/src/jit/helpers.hpp index 2abc300adeee..4525a7799171 100644 --- a/cpp/src/jit/helpers.hpp +++ b/cpp/src/jit/helpers.hpp @@ -38,7 +38,7 @@ std::map build_ptx_params(std::span ou template rmm::device_uvector to_device_vector(std::vector const& host, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { rmm::device_uvector device{host.size(), stream, mr}; @@ -50,7 +50,7 @@ template std::tuple>>, rmm::device_uvector> column_views_to_device(std::span views, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { std::vector>> handles; diff --git a/cpp/src/jit/row_ir.cpp b/cpp/src/jit/row_ir.cpp index b39a1d6f0e7b..816b8362eace 100644 --- a/cpp/src/jit/row_ir.cpp +++ b/cpp/src/jit/row_ir.cpp @@ -1000,7 +1000,7 @@ transform_args ast_converter::compute_table( table_view const& left_table, table_view const& right_table, std::string_view function_name, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { ast_converter converter{stream, mr, left_table, right_table}; @@ -1063,7 +1063,7 @@ transform_args ast_converter::filter(target target_id, table_view const& left_table, table_view const& right_table, std::string_view function_name, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto filter = ast::detail::predicate{expr}; diff --git a/cpp/src/jit/row_ir.hpp b/cpp/src/jit/row_ir.hpp index 2c506d91bfe3..f871866e5c1f 100644 --- a/cpp/src/jit/row_ir.hpp +++ b/cpp/src/jit/row_ir.hpp @@ -17,9 +17,10 @@ #include #include -#include #include +#include + #include #include #include @@ -116,8 +117,7 @@ struct [[nodiscard]] instance_context { std::vector input_vars_; ///< The input variables for the IR std::vector output_vars_; ///< The output variables for the IR std::unordered_multimap cse_nodes_; ///< multimap of IR nodes - rmm::cuda_stream_view - stream_; ///< The CUDA stream for any device operations during IR generation + cuda::stream_ref stream_; ///< The CUDA stream for any device operations during IR generation rmm::device_async_resource_ref mr_; ///< The device memory resource for any device memory allocation during IR generation @@ -125,7 +125,7 @@ struct [[nodiscard]] instance_context { friend struct ast_converter; friend struct node; - instance_context(rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + instance_context(cuda::stream_ref stream, rmm::device_async_resource_ref mr) : stream_(stream), mr_(mr) { } @@ -226,7 +226,7 @@ struct [[nodiscard]] instance_context { * @brief Get the CUDA stream for device operations during IR generation * @return The CUDA stream for device operations during IR generation */ - [[nodiscard]] rmm::cuda_stream_view get_stream() const { return stream_; } + [[nodiscard]] cuda::stream_ref get_stream() const { return stream_; } /** * @brief Get the device memory resource for device memory allocation during IR generation @@ -470,8 +470,7 @@ struct [[nodiscard]] node { struct [[nodiscard]] ast_converter { private: std::vector> output_irs_; ///< The output IR nodes - rmm::cuda_stream_view - stream_; ///< CUDA stream used for device memory operations and kernel launches. + cuda::stream_ref stream_; ///< CUDA stream used for device memory operations and kernel launches. rmm::device_async_resource_ref mr_; ///< Device memory resource used to allocate the returned table's device memory instance_context instance_; ///< The instance context used during the IR generation @@ -486,7 +485,7 @@ struct [[nodiscard]] ast_converter { * @param left_table Left input table referenced by expressions * @param right_table Right input table referenced by expressions */ - ast_converter(rmm::cuda_stream_view stream, + ast_converter(cuda::stream_ref stream, rmm::device_async_resource_ref mr, table_view left_table, table_view right_table) @@ -550,7 +549,7 @@ struct [[nodiscard]] ast_converter { table_view const& left_table, table_view const& right_table, std::string_view function_name, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -569,7 +568,7 @@ struct [[nodiscard]] ast_converter { table_view const& left_table, table_view const& right_table, std::string_view function_name, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); }; diff --git a/cpp/src/json/json_path.cu b/cpp/src/json/json_path.cu index a606f2996232..ff74c2167dde 100644 --- a/cpp/src/json/json_path.cu +++ b/cpp/src/json/json_path.cu @@ -649,7 +649,7 @@ class path_state : private parser { * @returns A pair containing the command buffer, and maximum stack depth required. */ std::pair>, int> build_command_buffer( - cudf::string_scalar const& json_path, rmm::cuda_stream_view stream) + cudf::string_scalar const& json_path, cuda::stream_ref stream) { std::string h_json_path = json_path.to_string(stream); path_state p_state(h_json_path.data(), static_cast(h_json_path.size())); @@ -969,7 +969,7 @@ __launch_bounds__(block_size) CUDF_KERNEL std::unique_ptr get_json_object(cudf::strings_column_view const& col, cudf::string_scalar const& json_path, get_json_object_options options, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // preprocess the json_path into a command buffer @@ -1006,7 +1006,7 @@ std::unique_ptr get_json_object(cudf::strings_column_view const& c auto cdv = column_device_view::create(col.parent(), stream); // preprocess sizes (returned in the offsets buffer) get_json_object_kernel - <<>>( + <<>>( *cdv, std::get<0>(preprocess).value().data(), sizes.data(), @@ -1034,7 +1034,7 @@ std::unique_ptr get_json_object(cudf::strings_column_view const& c 0, stream, cudf::get_current_device_resource_ref()}; get_json_object_kernel - <<>>( + <<>>( *cdv, std::get<0>(preprocess).value().data(), sizes.data(), @@ -1063,7 +1063,7 @@ std::unique_ptr get_json_object(cudf::strings_column_view const& c std::unique_ptr get_json_object(cudf::strings_column_view const& col, cudf::string_scalar const& json_path, get_json_object_options options, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/labeling/label_bins.cu b/cpp/src/labeling/label_bins.cu index 4d7fcfd9dbdb..d76569ad4b32 100644 --- a/cpp/src/labeling/label_bins.cu +++ b/cpp/src/labeling/label_bins.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -21,12 +21,12 @@ #include #include -#include #include #include #include #include +#include #include #include #include @@ -121,7 +121,7 @@ template std::unique_ptr label_bins(column_view const& input, column_view const& left_edges, column_view const& right_edges, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto output = make_numeric_column( @@ -179,7 +179,7 @@ struct bin_type_dispatcher { inclusive left_inclusive, column_view const& right_edges, inclusive right_inclusive, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(detail::is_supported_bin_type()) { @@ -200,7 +200,7 @@ std::unique_ptr label_bins(column_view const& input, inclusive left_inclusive, column_view const& right_edges, inclusive right_inclusive, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -235,7 +235,7 @@ std::unique_ptr label_bins(column_view const& input, inclusive left_inclusive, column_view const& right_edges, inclusive right_inclusive, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/merge/merge.cu b/cpp/src/merge/merge.cu index 4ba61ac6b52a..e52b22b0f38c 100644 --- a/cpp/src/merge/merge.cu +++ b/cpp/src/merge/merge.cu @@ -28,13 +28,13 @@ #include #include -#include #include #include #include #include #include +#include #include #include #include @@ -153,7 +153,7 @@ void materialize_bitmask(column_view const& left_col, bitmask_type* out_validity, size_type num_elements, index_type const* merged_indices, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { constexpr size_type BLOCK_SIZE{256}; detail::grid_1d grid_config{num_elements, BLOCK_SIZE}; @@ -167,19 +167,19 @@ void materialize_bitmask(column_view const& left_col, if (left_col.has_nulls()) { if (right_col.has_nulls()) { materialize_merged_bitmask_kernel - <<>>( + <<>>( left_valid, right_valid, out_validity, num_elements, merged_indices); CUDF_CUDA_TRY(cudaGetLastError()); } else { materialize_merged_bitmask_kernel - <<>>( + <<>>( left_valid, right_valid, out_validity, num_elements, merged_indices); CUDF_CUDA_TRY(cudaGetLastError()); } } else { if (right_col.has_nulls()) { materialize_merged_bitmask_kernel - <<>>( + <<>>( left_valid, right_valid, out_validity, num_elements, merged_indices); CUDF_CUDA_TRY(cudaGetLastError()); } else { @@ -187,7 +187,7 @@ void materialize_bitmask(column_view const& left_col, } } - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); } struct side_index_generator { @@ -218,7 +218,7 @@ index_vector generate_merged_indices(table_view const& left_table, std::vector const& column_order, std::vector const& null_precedence, bool nullable, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { size_type const left_size = left_table.num_rows(); size_type const right_size = right_table.num_rows(); @@ -275,7 +275,7 @@ index_vector generate_merged_indices(table_view const& left_table, ineq_op); } - stream.synchronize(); + CUDF_CHECK_CUDA(stream.get()); return merged_indices; } @@ -284,7 +284,7 @@ index_vector generate_merged_indices_nested(table_view const& left_table, std::vector const& column_order, std::vector const& null_precedence, bool nullable, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { size_type const left_size = left_table.num_rows(); size_type const right_size = right_table.num_rows(); @@ -339,7 +339,7 @@ struct column_merger { template ())> std::unique_ptr operator()(column_view const&, column_view const&, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref) const { CUDF_FAIL("Unsupported type for merge."); @@ -350,7 +350,7 @@ struct column_merger { template std::unique_ptr operator()(column_view const& lcol, column_view const& rcol, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const requires(is_rep_layout_compatible()) { @@ -423,7 +423,7 @@ template <> std::unique_ptr column_merger::operator()( column_view const& lcol, column_view const& rcol, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { return strings::detail::merge( @@ -435,7 +435,7 @@ template <> std::unique_ptr column_merger::operator()( column_view const& lcol, column_view const& rcol, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { auto result = cudf::dictionary::detail::merge( @@ -455,7 +455,7 @@ template <> std::unique_ptr column_merger::operator()( column_view const& lcol, column_view const& rcol, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { std::vector columns{lcol, rcol}; @@ -483,7 +483,7 @@ template <> std::unique_ptr column_merger::operator()( column_view const& lcol, column_view const& rcol, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { // merge each child. @@ -532,7 +532,7 @@ table_ptr_type merge(cudf::table_view const& left_table, std::vector const& key_cols, std::vector const& column_order, std::vector const& null_precedence, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // collect index columns for lhs, rhs, resp. @@ -602,7 +602,7 @@ table_ptr_type merge(std::vector const& tables_to_merge, std::vector const& key_cols, std::vector const& column_order, std::vector const& null_precedence, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (tables_to_merge.empty()) { return std::make_unique(); } @@ -685,7 +685,7 @@ std::unique_ptr merge(std::vector const& tables_to_merg std::vector const& key_cols, std::vector const& column_order, std::vector const& null_precedence, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/quantiles/quantile.cu b/cpp/src/quantiles/quantile.cu index 250bd81f41c8..01e9ea5e4110 100644 --- a/cpp/src/quantiles/quantile.cu +++ b/cpp/src/quantiles/quantile.cu @@ -21,12 +21,13 @@ #include #include -#include #include #include #include #include +#include +#include #include #include @@ -41,7 +42,7 @@ struct quantile_functor { std::vector const& q; interpolation interp; bool retain_types; - rmm::cuda_stream_view stream; + cuda::stream_ref stream; rmm::device_async_resource_ref mr; template @@ -132,7 +133,7 @@ std::unique_ptr quantile(column_view const& input, std::vector const& q, interpolation interp, bool retain_types, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto functor = quantile_functor{ @@ -150,7 +151,7 @@ std::unique_ptr quantile(column_view const& input, interpolation interp, column_view const& indices, bool exact, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (indices.is_empty()) { @@ -181,7 +182,7 @@ std::unique_ptr quantile(column_view const& input, interpolation interp, column_view const& ordered_indices, bool exact, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/quantiles/quantiles.cu b/cpp/src/quantiles/quantiles.cu index d3eac37020d1..22af344c4d52 100644 --- a/cpp/src/quantiles/quantiles.cu +++ b/cpp/src/quantiles/quantiles.cu @@ -17,10 +17,10 @@ #include #include -#include - #include #include +#include +#include #include #include @@ -33,7 +33,7 @@ std::unique_ptr
quantiles(table_view const& input, SortMapIterator sortmap, std::vector const& q, interpolation interp, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto quantile_idx_lookup = cuda::proclaim_return_type( @@ -61,7 +61,7 @@ std::unique_ptr
quantiles(table_view const& input, cudf::sorted is_input_sorted, std::vector const& column_order, std::vector const& null_precedence, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (q.empty()) { return empty_like(input); } @@ -90,7 +90,7 @@ std::unique_ptr
quantiles(table_view const& input, cudf::sorted is_input_sorted, std::vector const& column_order, std::vector const& null_precedence, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/quantiles/tdigest/tdigest.cu b/cpp/src/quantiles/tdigest/tdigest.cu index 73fee2faccbf..283f66bbaf80 100644 --- a/cpp/src/quantiles/tdigest/tdigest.cu +++ b/cpp/src/quantiles/tdigest/tdigest.cu @@ -19,13 +19,13 @@ #include #include -#include #include #include #include #include #include +#include #include #include #include @@ -174,7 +174,7 @@ CUDF_KERNEL void compute_percentiles_kernel(device_span tdigest_o */ std::unique_ptr compute_approx_percentiles(tdigest_column_view const& input, column_view const& percentiles, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { tdigest_column_view tdv(input); @@ -232,8 +232,8 @@ std::unique_ptr compute_approx_percentiles(tdigest_column_view const& in constexpr size_type block_size = 256; cudf::detail::grid_1d const grid(percentiles.size() * input.size(), block_size); - compute_percentiles_kernel<<>>( - {offsets.begin(), static_cast(offsets.size())}, + compute_percentiles_kernel<<>>( + {offsets.begin(), static_cast(offsets.size())}, *percentiles_cdv, centroids, tdv.min_begin(), @@ -251,7 +251,7 @@ std::unique_ptr make_tdigest_column(size_type num_rows, std::unique_ptr&& tdigest_offsets, std::unique_ptr&& min_values, std::unique_ptr&& max_values, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(tdigest_offsets->size() == num_rows + 1, @@ -284,7 +284,7 @@ std::unique_ptr make_tdigest_column(size_type num_rows, } std::unique_ptr make_empty_tdigests_column(size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto offsets = cudf::make_fixed_width_column( @@ -327,7 +327,7 @@ std::unique_ptr make_empty_tdigests_column(size_type num_rows, * * @returns An empty tdigest scalar. */ -std::unique_ptr make_empty_tdigest_scalar(rmm::cuda_stream_view stream, +std::unique_ptr make_empty_tdigest_scalar(cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto contents = make_empty_tdigests_column(1, stream, mr)->release(); @@ -339,7 +339,7 @@ std::unique_ptr make_empty_tdigest_scalar(rmm::cuda_stream_view stream, std::unique_ptr percentile_approx(tdigest_column_view const& input, column_view const& percentiles, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { tdigest_column_view tdv(input); @@ -367,7 +367,7 @@ std::unique_ptr percentile_approx(tdigest_column_view const& input, cudf::make_empty_column(type_id::FLOAT64), input.size(), cudf::detail::create_null_mask( - input.size(), mask_state::ALL_NULL, rmm::cuda_stream_view(stream), mr)); + input.size(), mask_state::ALL_NULL, cuda::stream_ref(stream), mr)); } // if any of the input digests are empty, nullify the corresponding output rows (values will be @@ -400,7 +400,7 @@ std::unique_ptr percentile_approx(tdigest_column_view const& input, std::unique_ptr percentile_approx(tdigest_column_view const& input, column_view const& percentiles, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/quantiles/tdigest/tdigest_aggregation.cu b/cpp/src/quantiles/tdigest/tdigest_aggregation.cu index a5ea03d1a8f8..ded66b749262 100644 --- a/cpp/src/quantiles/tdigest/tdigest_aggregation.cu +++ b/cpp/src/quantiles/tdigest/tdigest_aggregation.cu @@ -17,6 +17,7 @@ #include #include #include +#include #include #include #include @@ -382,7 +383,7 @@ CUDF_HOST_DEVICE constexpr inline double scale_func_k1(double quantile, // convert a single-row tdigest column to a scalar. std::unique_ptr to_tdigest_scalar(std::unique_ptr&& tdigest, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(tdigest->size() == 1, @@ -591,7 +592,7 @@ void generate_cluster_limits(int delta, size_type* group_num_clusters, size_type const* group_cluster_start, bool has_nulls, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); @@ -611,7 +612,7 @@ void generate_cluster_limits(int delta, constexpr size_type block_size = 256; cudf::detail::grid_1d const grid(num_gpu_groups, block_size); - generate_cluster_limits_kernel<<>>( + generate_cluster_limits_kernel<<>>( delta, num_gpu_groups, nearest_weight, @@ -647,7 +648,7 @@ template size_t compute_simple_cluster_count(int delta, GroupInfo group_info, cudf::device_span group_num_clusters, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const num_groups = group_num_clusters.size(); @@ -677,7 +678,7 @@ size_t compute_simple_cluster_count(int delta, * the terminology 'start' here instead of 'offsets' because our allocations strategy may * cause us to overallocate buffers within each group. */ -void compute_cluster_starts(cluster_info& cinfo, rmm::cuda_stream_view stream) +void compute_cluster_starts(cluster_info& cinfo, cuda::stream_ref stream) { auto const num_groups = cinfo.num_clusters.size(); auto cluster_size = cudf::detail::make_counting_transform_iterator( @@ -724,7 +725,7 @@ cluster_info generate_group_cluster_info(int delta, GroupInfo group_info, CumulativeWeight cumulative_weight, bool has_nulls, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -791,7 +792,7 @@ cluster_info generate_group_cluster_info(int delta, cinfo.cluster_wl = rmm::device_uvector(allocated_clusters, stream, temp_mr); // sync required after compute_cluster_starts() and before generate_cluster_limits() - stream.synchronize(); + cudf::detail::sync_stream(stream); // fill in the actual cluster weight limits. // if we are in the simple case, group_num_clusters will be updated here to reflect the accurate @@ -822,7 +823,7 @@ cluster_info generate_group_cluster_info(int delta, // cluster_start is returned as part of the output, so make sure to use the user supplied mr // instead of the current resource. cinfo.cluster_start = rmm::device_uvector(p_cluster_start, stream, mr); - stream.synchronize(); + cudf::detail::sync_stream(stream); } // if we are in the simple case we need to recompute the total clusters. allocated_cluster count @@ -835,7 +836,7 @@ cluster_info generate_group_cluster_info(int delta, cinfo.num_clusters.end()) : allocated_clusters; - stream.synchronize(); + cudf::detail::sync_stream(stream); return cinfo; } @@ -847,7 +848,7 @@ std::unique_ptr build_output_column(size_type num_rows, std::unique_ptr&& min_col, std::unique_ptr&& max_col, bool has_nulls, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // whether or not this weight is a stub @@ -1000,7 +1001,7 @@ std::unique_ptr compute_tdigests(int delta, std::unique_ptr&& max_col, cluster_info& cinfo, bool has_nulls, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // the output for each group is a column of data that represents the tdigest. since we want 1 row @@ -1051,7 +1052,7 @@ std::unique_ptr compute_tdigests(int delta, // Use `cub::DeviceReduce::ReduceByKey` instead of `thrust::reduce_by_key`: nvcc 13.0 // mis-compiles thrust's reduce-by-key kernel for sm_100, causing illegal memory accesses auto env = cuda::std::execution::env{ - cuda::std::execution::prop{cuda::get_stream_t{}, cuda::stream_ref{stream.value()}}, + cuda::std::execution::prop{cuda::get_stream_t{}, cuda::stream_ref{stream.get()}}, cuda::std::execution::prop{cuda::mr::get_memory_resource_t{}, cudf::get_current_device_resource_ref()}}; CUDF_CUDA_TRY(cub::DeviceReduce::ReduceByKey(keys, // keys in @@ -1126,7 +1127,7 @@ struct typed_group_tdigest { cudf::device_span group_valid_counts, size_type num_groups, int delta, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(cudf::is_numeric() || cudf::is_fixed_point()) { @@ -1150,7 +1151,7 @@ struct typed_group_tdigest { col.null_count() > 0, stream, mr); - stream.synchronize(); + cudf::detail::sync_stream(stream); return ret; } return generate_group_cluster_info( @@ -1213,7 +1214,7 @@ struct typed_reduce_tdigest { template std::unique_ptr operator()(column_view const& col, int delta, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(cudf::is_numeric() || cudf::is_fixed_point()) { @@ -1349,7 +1350,7 @@ std::pair, rmm::device_uvector> generate_mer tdigest_column_view const& tdv, GroupOffsetIter group_offsets, size_type num_groups, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); @@ -1396,7 +1397,7 @@ std::pair, rmm::device_uvector> generate_mer num_groups, centroid_offsets, centroid_offsets + 1, - stream.value())); + stream.get())); rmm::device_buffer temp_mem(temp_size, stream, temp_mr); CUDF_CUDA_TRY(cub::DeviceSegmentedSort::SortPairs(temp_mem.data(), @@ -1409,7 +1410,7 @@ std::pair, rmm::device_uvector> generate_mer num_groups, centroid_offsets, centroid_offsets + 1, - stream.value())); + stream.get())); return {std::move(output_means), std::move(output_weights)}; } @@ -1442,7 +1443,7 @@ std::unique_ptr merge_tdigests(tdigest_column_view const& tdv, size_t num_group_labels, size_type num_groups, int max_centroids, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // generate min and max values @@ -1559,7 +1560,7 @@ std::unique_ptr merge_tdigests(tdigest_column_view const& tdv, _p_group_labels.begin()); cudf::device_span p_group_labels(_p_group_labels); - stream.synchronize(); + cudf::detail::sync_stream(stream); return generate_group_cluster_info( delta, num_groups, @@ -1623,7 +1624,7 @@ std::unique_ptr merge_tdigests(tdigest_column_view const& tdv, std::unique_ptr reduce_tdigest(column_view const& col, int max_centroids, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (col.size() == 0) { return cudf::tdigest::detail::make_empty_tdigest_scalar(stream, mr); } @@ -1646,7 +1647,7 @@ struct group_offsets_fn { std::unique_ptr reduce_merge_tdigest(column_view const& input, int max_centroids, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { tdigest_column_view tdv(input); @@ -1668,7 +1669,7 @@ std::unique_ptr group_tdigest(column_view const& col, cudf::device_span group_valid_counts, size_type num_groups, int max_centroids, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (col.size() == 0) { return cudf::tdigest::detail::make_empty_tdigests_column(1, stream, mr); } @@ -1691,7 +1692,7 @@ std::unique_ptr group_merge_tdigest(column_view const& input, cudf::device_span group_labels, size_type num_groups, int max_centroids, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { tdigest_column_view tdv(input); diff --git a/cpp/src/replace/clamp.cu b/cpp/src/replace/clamp.cu index 836af9b5a86a..7830efd511f8 100644 --- a/cpp/src/replace/clamp.cu +++ b/cpp/src/replace/clamp.cu @@ -27,12 +27,12 @@ #include #include -#include #include #include #include #include +#include #include #include @@ -77,7 +77,7 @@ std::unique_ptr clamp_string_column(strings_column_view const& inp ReplaceScalarIterator lo_replace_itr, OptionalScalarIterator hi_itr, ReplaceScalarIterator hi_replace_itr, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto input_device_column = column_device_view::create(input.parent(), stream); @@ -121,7 +121,7 @@ std::unique_ptr clamp_dictionary_column(dictionary_column_view con scalar const& lo_replace, scalar const& hi, scalar const& hi_replace, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // add lo_replace and hi_replace to keys @@ -190,7 +190,7 @@ std::unique_ptr clamper(column_view const& input, ReplaceScalarIterator lo_replace_itr, OptionalScalarIterator hi_itr, ReplaceScalarIterator hi_replace_itr, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(cudf::is_fixed_width()) { @@ -240,7 +240,7 @@ std::unique_ptr clamper(column_view const& input, ReplaceScalarIterator lo_replace_itr, OptionalScalarIterator hi_itr, ReplaceScalarIterator hi_replace_itr, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(std::is_same_v) { @@ -255,7 +255,7 @@ std::unique_ptr clamp(column_view const& input, ReplaceScalarIterator lo_replace_itr, OptionalScalarIterator hi_itr, ReplaceScalarIterator hi_replace_itr, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return clamper(input, lo_itr, lo_replace_itr, hi_itr, hi_replace_itr, stream, mr); @@ -268,7 +268,7 @@ struct dispatch_clamp { scalar const& lo_replace, scalar const& hi, scalar const& hi_replace, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(cudf::have_same_types(input, lo), @@ -296,7 +296,7 @@ std::unique_ptr dispatch_clamp::operator()( scalar const& lo_replace, scalar const& hi, scalar const& hi_replace, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FAIL("clamp for list_view not supported"); @@ -308,7 +308,7 @@ std::unique_ptr dispatch_clamp::operator()(column_view cons scalar const& lo_replace, scalar const& hi, scalar const& hi_replace, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FAIL("clamp for struct_view not supported"); @@ -320,7 +320,7 @@ std::unique_ptr dispatch_clamp::operator()(column_view con scalar const&, scalar const&, scalar const&, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref) { CUDF_UNREACHABLE("clamp type-dispatch error"); @@ -341,7 +341,7 @@ std::unique_ptr clamp(column_view const& input, scalar const& lo_replace, scalar const& hi, scalar const& hi_replace, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS( @@ -379,7 +379,7 @@ std::unique_ptr clamp(column_view const& input, scalar const& lo_replace, scalar const& hi, scalar const& hi_replace, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -390,7 +390,7 @@ std::unique_ptr clamp(column_view const& input, std::unique_ptr clamp(column_view const& input, scalar const& lo, scalar const& hi, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/replace/nans.cu b/cpp/src/replace/nans.cu index c943b7a78be8..23b250a6ad74 100644 --- a/cpp/src/replace/nans.cu +++ b/cpp/src/replace/nans.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include @@ -16,11 +16,11 @@ #include #include -#include #include #include #include +#include #include #include @@ -33,7 +33,7 @@ struct replace_nans_functor { std::unique_ptr operator()(column_view const& input, Replacement const& replacement, bool replacement_nullable, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(std::is_floating_point_v) { @@ -75,7 +75,7 @@ struct replace_nans_functor { std::unique_ptr replace_nans(column_view const& input, column_view const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input.size() == replacement.size(), @@ -92,7 +92,7 @@ std::unique_ptr replace_nans(column_view const& input, std::unique_ptr replace_nans(column_view const& input, scalar const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return type_dispatcher( @@ -103,7 +103,7 @@ std::unique_ptr replace_nans(column_view const& input, std::unique_ptr replace_nans(column_view const& input, column_view const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -112,7 +112,7 @@ std::unique_ptr replace_nans(column_view const& input, std::unique_ptr replace_nans(column_view const& input, scalar const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -144,7 +144,7 @@ struct normalize_nans_and_zeros_kernel_forwarder { template void operator()(cudf::column_device_view in, cudf::mutable_column_device_view out, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) requires(std::is_floating_point_v) { thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), @@ -167,7 +167,7 @@ struct normalize_nans_and_zeros_kernel_forwarder { namespace cudf { namespace detail { -void normalize_nans_and_zeros(mutable_column_view in_out, rmm::cuda_stream_view stream) +void normalize_nans_and_zeros(mutable_column_view in_out, cuda::stream_ref stream) { if (in_out.is_empty()) { return; } CUDF_EXPECTS( @@ -190,7 +190,7 @@ void normalize_nans_and_zeros(mutable_column_view in_out, rmm::cuda_stream_view } std::unique_ptr normalize_nans_and_zeros(column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // output. copies the input @@ -217,7 +217,7 @@ std::unique_ptr normalize_nans_and_zeros(column_view const& input, * @param mr Device memory resource used to allocate the returned column's device memory. */ std::unique_ptr normalize_nans_and_zeros(column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -235,7 +235,7 @@ std::unique_ptr normalize_nans_and_zeros(column_view const& input, * @param[in, out] in_out mutable_column_view representing input data. data is processed in-place * @param stream CUDA stream used for device memory operations and kernel launches */ -void normalize_nans_and_zeros(mutable_column_view& in_out, rmm::cuda_stream_view stream) +void normalize_nans_and_zeros(mutable_column_view& in_out, cuda::stream_ref stream) { CUDF_FUNC_RANGE(); detail::normalize_nans_and_zeros(in_out, stream); diff --git a/cpp/src/replace/nulls.cu b/cpp/src/replace/nulls.cu index 90c91af54d23..b1dc4de69a7a 100644 --- a/cpp/src/replace/nulls.cu +++ b/cpp/src/replace/nulls.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -32,11 +32,11 @@ #include #include -#include #include #include #include +#include #include #include @@ -100,7 +100,7 @@ struct replace_nulls_column_kernel_forwarder { template ())> std::unique_ptr operator()(cudf::column_view const& input, cudf::column_view const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { cudf::size_type nrows = input.size(); @@ -127,7 +127,7 @@ struct replace_nulls_column_kernel_forwarder { 0, stream, cudf::get_current_device_resource_ref()); cudf::size_type* valid_count = valid_counter.data(); - replace<<>>( + replace<<>>( *device_in, *device_replacement, *device_out, valid_count); CUDF_CUDA_TRY(cudaGetLastError()); @@ -141,7 +141,7 @@ struct replace_nulls_column_kernel_forwarder { template ())> std::unique_ptr operator()(cudf::column_view const&, cudf::column_view const&, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref) { CUDF_FAIL("No specialization exists for the given type."); @@ -152,7 +152,7 @@ template <> std::unique_ptr replace_nulls_column_kernel_forwarder::operator()( cudf::column_view const& input, cudf::column_view const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto d_input = cudf::column_device_view::create(input, stream); @@ -179,7 +179,7 @@ template <> std::unique_ptr replace_nulls_column_kernel_forwarder::operator()( cudf::column_view const& input, cudf::column_view const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { cudf::dictionary_column_view dict_input(input); @@ -202,7 +202,7 @@ struct replace_nulls_scalar_kernel_forwarder { template ()>* = nullptr> std::unique_ptr operator()(cudf::column_view const& input, cudf::scalar const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS( @@ -228,7 +228,7 @@ struct replace_nulls_scalar_kernel_forwarder { template ()>* = nullptr> std::unique_ptr operator()(cudf::column_view const&, cudf::scalar const&, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref) { CUDF_FAIL("No specialization exists for the given type."); @@ -239,7 +239,7 @@ template <> std::unique_ptr replace_nulls_scalar_kernel_forwarder::operator()( cudf::column_view const& input, cudf::scalar const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS( @@ -253,7 +253,7 @@ template <> std::unique_ptr replace_nulls_scalar_kernel_forwarder::operator()( cudf::column_view const& input, cudf::scalar const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { cudf::dictionary_column_view dict_input(input); @@ -266,7 +266,7 @@ std::unique_ptr replace_nulls_scalar_kernel_forwarder::operator()< std::unique_ptr replace_nulls_policy_impl(cudf::column_view const& input, cudf::replace_policy const& replace_policy, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto device_in = cudf::column_device_view::create(input, stream); @@ -312,7 +312,7 @@ namespace detail { std::unique_ptr replace_nulls(cudf::column_view const& input, cudf::column_view const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS( @@ -328,7 +328,7 @@ std::unique_ptr replace_nulls(cudf::column_view const& input, std::unique_ptr replace_nulls(cudf::column_view const& input, cudf::scalar const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.is_empty()) { return cudf::empty_like(input); } @@ -342,7 +342,7 @@ std::unique_ptr replace_nulls(cudf::column_view const& input, std::unique_ptr replace_nulls(cudf::column_view const& input, cudf::replace_policy const& replace_policy, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.is_empty()) { return cudf::empty_like(input); } @@ -355,7 +355,7 @@ std::unique_ptr replace_nulls(cudf::column_view const& input, std::unique_ptr replace_nulls(cudf::column_view const& input, cudf::column_view const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -364,7 +364,7 @@ std::unique_ptr replace_nulls(cudf::column_view const& input, std::unique_ptr replace_nulls(cudf::column_view const& input, cudf::scalar const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -373,7 +373,7 @@ std::unique_ptr replace_nulls(cudf::column_view const& input, std::unique_ptr replace_nulls(column_view const& input, replace_policy const& replace_policy, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/replace/replace.cu b/cpp/src/replace/replace.cu index af16841d482b..a7f5d8f64171 100644 --- a/cpp/src/replace/replace.cu +++ b/cpp/src/replace/replace.cu @@ -3,7 +3,7 @@ * SPDX-FileCopyrightText: Copyright 2018 BlazingDB, Inc. * SPDX-FileCopyrightText: Copyright 2018 Cristhian Alberto Gonzales Castillo * SPDX-FileCopyrightText: Copyright 2018 Alexander Ocsa - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ // clang-format on @@ -48,10 +48,9 @@ #include #include -#include - #include #include +#include #include #include #include @@ -169,7 +168,7 @@ struct replace_kernel_forwarder { std::unique_ptr operator()(cudf::column_view const& input_col, cudf::column_view const& values_to_replace, cudf::column_view const& replacement_values, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { cudf::detail::device_scalar valid_counter( @@ -201,12 +200,12 @@ struct replace_kernel_forwarder { auto device_values_to_replace = cudf::column_device_view::create(values_to_replace, stream); auto device_replacement_values = cudf::column_device_view::create(replacement_values, stream); - replace<<>>(*device_in, - *device_out, - valid_count, - output_view.size(), - *device_values_to_replace, - *device_replacement_values); + replace<<>>(*device_in, + *device_out, + valid_count, + output_view.size(), + *device_values_to_replace, + *device_replacement_values); CUDF_CUDA_TRY(cudaGetLastError()); if (output_view.nullable()) { @@ -219,7 +218,7 @@ struct replace_kernel_forwarder { std::unique_ptr operator()(cudf::column_view const&, cudf::column_view const&, cudf::column_view const&, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref) { CUDF_FAIL("No specialization exists for this type"); @@ -231,7 +230,7 @@ std::unique_ptr replace_kernel_forwarder::operator() replace_kernel_forwarder::operator() find_and_replace_all(cudf::column_view const& input_col, cudf::column_view const& values_to_replace, cudf::column_view const& replacement_values, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(values_to_replace.size() == replacement_values.size(), @@ -328,7 +327,7 @@ std::unique_ptr find_and_replace_all(cudf::column_view const& inpu std::unique_ptr find_and_replace_all(cudf::column_view const& input_col, cudf::column_view const& values_to_replace, cudf::column_view const& replacement_values, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return detail::find_and_replace_all(input_col, values_to_replace, replacement_values, stream, mr); diff --git a/cpp/src/reshape/byte_cast.cu b/cpp/src/reshape/byte_cast.cu index 3e563a93cd1e..302d99502173 100644 --- a/cpp/src/reshape/byte_cast.cu +++ b/cpp/src/reshape/byte_cast.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -15,10 +15,10 @@ #include #include -#include #include #include +#include #include #include @@ -44,7 +44,7 @@ struct byte_list_conversion_dispatcher { template std::unique_ptr operator()(column_view const& input, flip_endianness configuration, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { return byte_list_conversion_fn::invoke(input, configuration, stream, mr); @@ -55,7 +55,7 @@ template struct byte_list_conversion_fn()>> { static std::unique_ptr invoke(column_view const& input, flip_endianness configuration, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.size() == 0) { return cudf::lists::detail::make_empty_lists_column(output_type); } @@ -111,7 +111,7 @@ template struct byte_list_conversion_fn>> { static std::unique_ptr invoke(column_view const& input, flip_endianness, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.size() == 0) { return cudf::lists::detail::make_empty_lists_column(output_type); } @@ -152,7 +152,7 @@ struct byte_list_conversion_fn byte_cast(column_view const& input, flip_endianness endian_configuration, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return type_dispatcher( @@ -163,7 +163,7 @@ std::unique_ptr byte_cast(column_view const& input, std::unique_ptr byte_cast(column_view const& input, flip_endianness endian_configuration, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/reshape/interleave_columns.cu b/cpp/src/reshape/interleave_columns.cu index 5867bbdc82e1..f372a8714ac6 100644 --- a/cpp/src/reshape/interleave_columns.cu +++ b/cpp/src/reshape/interleave_columns.cu @@ -20,11 +20,11 @@ #include #include -#include #include #include #include +#include #include #include @@ -45,7 +45,7 @@ struct interleave_columns_functor { template std::unique_ptr operator()(table_view const& input, bool create_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return interleave_columns_impl{}(input, create_mask, stream, mr); @@ -56,7 +56,7 @@ template struct interleave_columns_impl>> { std::unique_ptr operator()(table_view const& lists_columns, bool create_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return lists::detail::interleave_columns(lists_columns, create_mask, stream, mr); @@ -67,7 +67,7 @@ template struct interleave_columns_impl>> { std::unique_ptr operator()(table_view const& structs_columns, bool create_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // We can safely call `column(0)` as the number of columns is known to be non zero. @@ -152,7 +152,7 @@ template struct interleave_columns_impl>> { std::unique_ptr operator()(table_view const& strings_columns, bool, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto num_columns = strings_columns.num_columns(); @@ -184,7 +184,7 @@ template struct interleave_columns_impl()>> { std::unique_ptr operator()(table_view const& input, bool create_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto arch_column = input.column(0); @@ -234,7 +234,7 @@ struct interleave_columns_impl()>> { } // anonymous namespace std::unique_ptr interleave_columns(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input.num_columns() > 0, "input must have at least one column to determine dtype."); @@ -255,7 +255,7 @@ std::unique_ptr interleave_columns(table_view const& input, } // namespace detail std::unique_ptr interleave_columns(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/reshape/table_to_array.cu b/cpp/src/reshape/table_to_array.cu index 9aa3fbe7a91d..0bc46b15af9c 100644 --- a/cpp/src/reshape/table_to_array.cu +++ b/cpp/src/reshape/table_to_array.cu @@ -6,6 +6,7 @@ #include #include #include +#include #include #include #include @@ -15,12 +16,12 @@ #include #include -#include #include #include #include #include +#include #include #include @@ -31,7 +32,7 @@ namespace { template void table_to_array_impl(table_view const& input, device_span output, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const num_columns = input.num_columns(); auto const num_rows = input.num_rows(); @@ -65,14 +66,15 @@ void table_to_array_impl(table_view const& input, cuda::constant_iterator sizes(static_cast(item_size * num_rows)); cudf::detail::batched_memcpy_async( - d_srcs.begin(), d_dsts.begin(), sizes, num_columns, stream.value()); - stream.synchronize(); // ensures h_srcs and h_dsts are not destroyed before the copy is done + d_srcs.begin(), d_dsts.begin(), sizes, num_columns, stream.get()); + cudf::detail::sync_stream( + stream); // ensures h_srcs and h_dsts are not destroyed before the copy is done } struct table_to_array_dispatcher { table_view const& input; device_span output; - rmm::cuda_stream_view stream; + cuda::stream_ref stream; template ())> void operator()() const @@ -91,7 +93,7 @@ struct table_to_array_dispatcher { void table_to_array(table_view const& input, device_span output, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (input.num_columns() == 0) return; @@ -105,7 +107,7 @@ void table_to_array(table_view const& input, void table_to_array(table_view const& input, device_span output, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); cudf::detail::table_to_array(input, output, stream); diff --git a/cpp/src/reshape/tile.cu b/cpp/src/reshape/tile.cu index 4ce819e839f7..31153cbb0ccf 100644 --- a/cpp/src/reshape/tile.cu +++ b/cpp/src/reshape/tile.cu @@ -15,9 +15,8 @@ #include #include -#include - -#include +#include +#include #include @@ -33,7 +32,7 @@ struct tile_functor { namespace detail { std::unique_ptr
tile(table_view const& in, size_type count, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(count >= 0, "Count cannot be negative"); @@ -52,7 +51,7 @@ std::unique_ptr
tile(table_view const& in, std::unique_ptr
tile(table_view const& in, size_type count, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/round/round.cu b/cpp/src/round/round.cu index 85bb9160bfcf..6f502009e09f 100644 --- a/cpp/src/round/round.cu +++ b/cpp/src/round/round.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -18,9 +18,9 @@ #include #include -#include #include +#include #include #include @@ -211,7 +211,7 @@ struct half_even_fixed_point { template typename RoundFunctor> std::unique_ptr round_with(column_view const& input, int32_t decimal_places, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(not cudf::is_fixed_point()) { @@ -244,7 +244,7 @@ std::unique_ptr round_with(column_view const& input, template typename RoundFunctor> std::unique_ptr round_with(column_view const& input, int32_t decimal_places, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(cudf::is_fixed_point()) { @@ -310,7 +310,7 @@ struct round_type_dispatcher { std::unique_ptr operator()(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(is_supported_round_type()) { @@ -350,7 +350,7 @@ struct round_dispatch_fn { std::unique_ptr operator()(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(is_supported()) { @@ -377,7 +377,7 @@ struct round_dispatch_fn { std::unique_ptr round(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) { CUDF_EXPECTS(cudf::is_numeric(input.type()) || cudf::is_fixed_point(input.type()), @@ -398,7 +398,7 @@ std::unique_ptr round(column_view const& input, std::unique_ptr round_decimal(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) { CUDF_EXPECTS(cudf::is_integral_not_bool(input.type()) || cudf::is_fixed_point(input.type()), @@ -422,7 +422,7 @@ std::unique_ptr round_decimal(column_view const& input, std::unique_ptr round(column_view const& input, int32_t decimal_places, rounding_method method, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -432,7 +432,7 @@ std::unique_ptr round(column_view const& input, std::unique_ptr round_decimal(column_view const& input, int32_t decimal_places, rounding_method method, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/row_operator/row_operators.cu b/cpp/src/row_operator/row_operators.cu index 6f91e3d076df..69d6f1d089ab 100644 --- a/cpp/src/row_operator/row_operators.cu +++ b/cpp/src/row_operator/row_operators.cu @@ -14,6 +14,7 @@ #include #include #include +#include #include #include #include @@ -296,7 +297,7 @@ auto decompose_structs(table_view table, * This helper function generates dremel data for any list-type columns in a * table. This data is necessary for lexicographic comparisons. */ -auto list_lex_preprocess(table_view const& table, rmm::cuda_stream_view stream) +auto list_lex_preprocess(table_view const& table, cuda::stream_ref stream) { std::vector dremel_data; auto const num_list_columns = std::count_if( @@ -411,7 +412,7 @@ namespace { auto replace_child(column_view const& input, column_view const& new_child, std::vector>& out_cols, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const make_output = [&input](auto const& offsets_cv, auto const& child_cv) { @@ -460,7 +461,7 @@ auto replace_child(column_view const& input, */ auto compute_ranks(column_view const& input, null_order column_null_order, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return cudf::detail::rank(input, @@ -493,7 +494,7 @@ auto compute_ranks(column_view const& input, std::pair>> transform_lists_of_structs( column_view const& input, null_order column_null_order, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { std::vector> out_cols; @@ -560,7 +561,7 @@ std::tuple> out_cols_lhs; @@ -637,7 +638,7 @@ std::shared_ptr preprocessed_table::create( host_span column_order, host_span null_precedence, bool has_ranked_children, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { check_lex_compatibility(preprocessed_input); @@ -648,7 +649,7 @@ std::shared_ptr preprocessed_table::create( null_precedence, stream, cudf::get_current_device_resource_ref()); auto d_depths = detail::make_device_uvector_async( verticalized_col_depths, stream, cudf::get_current_device_resource_ref()); - stream.synchronize(); + cudf::detail::sync_stream(stream); if (detail::has_nested_columns(preprocessed_input)) { auto [dremel_data, d_dremel_device_view] = list_lex_preprocess(preprocessed_input, stream); @@ -676,7 +677,7 @@ std::shared_ptr preprocessed_table::create( table_view const& input, host_span column_order, host_span null_precedence, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto [decomposed_input, new_column_order, new_null_precedence, verticalized_col_depths] = decompose_structs(input, decompose_lists_column::NO, column_order, null_precedence); @@ -720,7 +721,7 @@ preprocessed_table::create(table_view const& lhs, table_view const& rhs, host_span column_order, host_span null_precedence, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { check_shape_compatibility(lhs, rhs); @@ -833,7 +834,7 @@ two_table_comparator::two_table_comparator(table_view const& left, table_view const& right, host_span column_order, host_span null_precedence, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { std::tie(d_left_table, d_right_table) = preprocessed_table::create(left, right, column_order, null_precedence, stream); @@ -844,7 +845,7 @@ two_table_comparator::two_table_comparator(table_view const& left, namespace equality { std::shared_ptr preprocessed_table::create( - table_view const& t, rmm::cuda_stream_view stream, rmm::device_async_resource_ref temp_mr) + table_view const& t, cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) { check_eq_compatibility(t); @@ -860,7 +861,7 @@ std::shared_ptr preprocessed_table::create( two_table_comparator::two_table_comparator(table_view const& left, table_view const& right, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) : d_left_table{preprocessed_table::create(left, stream, temp_mr)}, d_right_table{preprocessed_table::create(right, stream, temp_mr)} diff --git a/cpp/src/stream_compaction/apply_boolean_mask.cu b/cpp/src/stream_compaction/apply_boolean_mask.cu index 41e6d9f79fad..ca3e414aaee7 100644 --- a/cpp/src/stream_compaction/apply_boolean_mask.cu +++ b/cpp/src/stream_compaction/apply_boolean_mask.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -16,7 +16,7 @@ #include #include -#include +#include #include @@ -64,7 +64,7 @@ namespace detail { std::unique_ptr
apply_mask(table_view const& input, column_view const& boolean_mask, mask_type mask_kind, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (boolean_mask.is_empty()) { return empty_like(input); } @@ -98,7 +98,7 @@ std::unique_ptr
apply_mask(table_view const& input, */ std::unique_ptr
apply_boolean_mask(table_view const& input, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -107,7 +107,7 @@ std::unique_ptr
apply_boolean_mask(table_view const& input, std::unique_ptr
apply_deletion_mask(table_view const& input, column_view const& deletion_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/stream_compaction/distinct.cu b/cpp/src/stream_compaction/distinct.cu index 70dc44b027d0..07a348b12c4a 100644 --- a/cpp/src/stream_compaction/distinct.cu +++ b/cpp/src/stream_compaction/distinct.cu @@ -21,13 +21,13 @@ #include #include -#include #include #include #include #include #include +#include #include #include @@ -77,7 +77,7 @@ rmm::device_uvector distinct_indices(table_view const& input, duplicate_keep_option keep, null_equality nulls_equal, nan_equality nans_equal, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const num_rows = input.num_rows(); @@ -107,7 +107,7 @@ rmm::device_uvector distinct_indices(table_view const& input, {}, {}, rmm::mr::polymorphic_allocator{temp_mr}, - stream.value()}; + stream.get()}; return reduce_func(set); }; @@ -147,7 +147,7 @@ std::unique_ptr
distinct(table_view const& input, duplicate_keep_option keep, null_equality nulls_equal, nan_equality nans_equal, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.num_rows() == 0 or input.num_columns() == 0 or keys.empty()) { @@ -175,7 +175,7 @@ std::unique_ptr
distinct(table_view const& input, duplicate_keep_option keep, null_equality nulls_equal, nan_equality nans_equal, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -186,7 +186,7 @@ std::unique_ptr distinct_indices(table_view const& input, duplicate_keep_option keep, null_equality nulls_equal, nan_equality nans_equal, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/stream_compaction/distinct_helpers.cu b/cpp/src/stream_compaction/distinct_helpers.cu index 2596496c8774..afd290d8c254 100644 --- a/cpp/src/stream_compaction/distinct_helpers.cu +++ b/cpp/src/stream_compaction/distinct_helpers.cu @@ -10,13 +10,13 @@ #include #include -#include #include #include #include #include #include +#include #include namespace cudf::detail { @@ -24,7 +24,7 @@ namespace cudf::detail { void initialize_reduction_results(size_type* results, size_type num_rows, duplicate_keep_option keep, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { thrust::uninitialized_fill( rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), @@ -37,7 +37,7 @@ size_type copy_reduction_results(size_type const* results, size_type num_rows, size_type* output, duplicate_keep_option keep, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const output_end = [&] { if (keep == duplicate_keep_option::KEEP_NONE) { diff --git a/cpp/src/stream_compaction/distinct_helpers.cuh b/cpp/src/stream_compaction/distinct_helpers.cuh index 7649ea081380..d0c97953c619 100644 --- a/cpp/src/stream_compaction/distinct_helpers.cuh +++ b/cpp/src/stream_compaction/distinct_helpers.cuh @@ -11,7 +11,6 @@ #include #include -#include #include #include #include @@ -20,6 +19,7 @@ #include #include #include +#include #include namespace cudf::detail { @@ -27,14 +27,14 @@ namespace cudf::detail { template rmm::device_uvector reduce_by_row_keep_any(Set& set, size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto output_indices = rmm::device_uvector(num_rows, stream, mr); auto const iter = cuda::counting_iterator{0}; - set.insert_async(iter, iter + num_rows, stream.value()); - auto const output_end = set.retrieve_all(output_indices.begin(), stream.value()); + set.insert_async(iter, iter + num_rows, stream.get()); + auto const output_end = set.retrieve_all(output_indices.begin(), stream.get()); output_indices.resize(cuda::std::distance(output_indices.begin(), output_end), stream); return output_indices; } @@ -43,7 +43,7 @@ template rmm::device_uvector reduce_by_row_keep_first_last_none(Set& set, size_type num_rows, duplicate_keep_option keep, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto output_indices = rmm::device_uvector(num_rows, stream, mr); diff --git a/cpp/src/stream_compaction/distinct_helpers.hpp b/cpp/src/stream_compaction/distinct_helpers.hpp index c7903aae7672..93ee8da0634a 100644 --- a/cpp/src/stream_compaction/distinct_helpers.hpp +++ b/cpp/src/stream_compaction/distinct_helpers.hpp @@ -14,7 +14,6 @@ #include #include -#include #include #include #include @@ -24,6 +23,7 @@ #include #include #include +#include #include #include @@ -49,13 +49,13 @@ auto constexpr reduction_init_value(duplicate_keep_option keep) CUDF_HIDDEN void initialize_reduction_results(size_type* results, size_type num_rows, duplicate_keep_option keep, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); CUDF_HIDDEN size_type copy_reduction_results(size_type const* results, size_type num_rows, size_type* output, duplicate_keep_option keep, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); struct distinct_precomputed_hash { CUDF_HOST_DEVICE constexpr distinct_precomputed_hash(hash_value_type const* hashes) @@ -96,7 +96,7 @@ using distinct_set_t = cuco::static_set rmm::device_uvector reduce_by_row_keep_any(Set& set, size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -118,7 +118,7 @@ rmm::device_uvector reduce_by_row_keep_first_last_none( Set& set, size_type num_rows, duplicate_keep_option keep, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -136,7 +136,7 @@ template rmm::device_uvector reduce_by_row(Set& set, size_type num_rows, duplicate_keep_option keep, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (keep == duplicate_keep_option::KEEP_ANY) { diff --git a/cpp/src/stream_compaction/distinct_helpers_flat_nan_equal_any.cu b/cpp/src/stream_compaction/distinct_helpers_flat_nan_equal_any.cu index 91c112e03595..85bc4d3626e7 100644 --- a/cpp/src/stream_compaction/distinct_helpers_flat_nan_equal_any.cu +++ b/cpp/src/stream_compaction/distinct_helpers_flat_nan_equal_any.cu @@ -9,10 +9,11 @@ #include #include -#include #include #include +#include + namespace cudf::detail { template rmm::device_uvector reduce_by_row_keep_any( @@ -21,7 +22,7 @@ template rmm::device_uvector reduce_by_row_keep_any( cudf::nullate::DYNAMIC, cudf::detail::row::equality::nan_equal_physical_equality_comparator>>& set, size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace cudf::detail diff --git a/cpp/src/stream_compaction/distinct_helpers_flat_nan_equal_ordered.cu b/cpp/src/stream_compaction/distinct_helpers_flat_nan_equal_ordered.cu index f266862dcce6..fab477fc1d19 100644 --- a/cpp/src/stream_compaction/distinct_helpers_flat_nan_equal_ordered.cu +++ b/cpp/src/stream_compaction/distinct_helpers_flat_nan_equal_ordered.cu @@ -10,10 +10,11 @@ #include #include -#include #include #include +#include + namespace cudf::detail { template rmm::device_uvector reduce_by_row_keep_first_last_none( @@ -23,7 +24,7 @@ template rmm::device_uvector reduce_by_row_keep_first_last_none( cudf::detail::row::equality::nan_equal_physical_equality_comparator>>& set, size_type num_rows, duplicate_keep_option keep, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace cudf::detail diff --git a/cpp/src/stream_compaction/distinct_helpers_flat_nan_unequal_any.cu b/cpp/src/stream_compaction/distinct_helpers_flat_nan_unequal_any.cu index ce224e6d91d0..6a9ec3489b52 100644 --- a/cpp/src/stream_compaction/distinct_helpers_flat_nan_unequal_any.cu +++ b/cpp/src/stream_compaction/distinct_helpers_flat_nan_unequal_any.cu @@ -9,10 +9,11 @@ #include #include -#include #include #include +#include + namespace cudf::detail { template rmm::device_uvector reduce_by_row_keep_any( @@ -21,7 +22,7 @@ template rmm::device_uvector reduce_by_row_keep_any( cudf::nullate::DYNAMIC, cudf::detail::row::equality::physical_equality_comparator>>& set, size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace cudf::detail diff --git a/cpp/src/stream_compaction/distinct_helpers_flat_nan_unequal_ordered.cu b/cpp/src/stream_compaction/distinct_helpers_flat_nan_unequal_ordered.cu index 40a9a327e6c4..cf85ac5bc74d 100644 --- a/cpp/src/stream_compaction/distinct_helpers_flat_nan_unequal_ordered.cu +++ b/cpp/src/stream_compaction/distinct_helpers_flat_nan_unequal_ordered.cu @@ -10,10 +10,11 @@ #include #include -#include #include #include +#include + namespace cudf::detail { template rmm::device_uvector reduce_by_row_keep_first_last_none( @@ -23,7 +24,7 @@ template rmm::device_uvector reduce_by_row_keep_first_last_none( cudf::detail::row::equality::physical_equality_comparator>>& set, size_type num_rows, duplicate_keep_option keep, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace cudf::detail diff --git a/cpp/src/stream_compaction/distinct_helpers_nested_nan_equal_any.cu b/cpp/src/stream_compaction/distinct_helpers_nested_nan_equal_any.cu index e60ad6a2547e..59a60f21d198 100644 --- a/cpp/src/stream_compaction/distinct_helpers_nested_nan_equal_any.cu +++ b/cpp/src/stream_compaction/distinct_helpers_nested_nan_equal_any.cu @@ -9,10 +9,11 @@ #include #include -#include #include #include +#include + namespace cudf::detail { template rmm::device_uvector reduce_by_row_keep_any( @@ -22,7 +23,7 @@ template rmm::device_uvector reduce_by_row_keep_any( cudf::detail::row::equality::nan_equal_physical_equality_comparator>, distinct_precomputed_hash>& set, size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace cudf::detail diff --git a/cpp/src/stream_compaction/distinct_helpers_nested_nan_equal_ordered.cu b/cpp/src/stream_compaction/distinct_helpers_nested_nan_equal_ordered.cu index 3030a4c63e99..bf97e75e50fb 100644 --- a/cpp/src/stream_compaction/distinct_helpers_nested_nan_equal_ordered.cu +++ b/cpp/src/stream_compaction/distinct_helpers_nested_nan_equal_ordered.cu @@ -10,10 +10,11 @@ #include #include -#include #include #include +#include + namespace cudf::detail { template rmm::device_uvector reduce_by_row_keep_first_last_none( @@ -23,7 +24,7 @@ template rmm::device_uvector reduce_by_row_keep_first_last_none( cudf::detail::row::equality::nan_equal_physical_equality_comparator>>& set, size_type num_rows, duplicate_keep_option keep, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace cudf::detail diff --git a/cpp/src/stream_compaction/distinct_helpers_nested_nan_unequal_any.cu b/cpp/src/stream_compaction/distinct_helpers_nested_nan_unequal_any.cu index 111ee665bd53..2277d1b98c1e 100644 --- a/cpp/src/stream_compaction/distinct_helpers_nested_nan_unequal_any.cu +++ b/cpp/src/stream_compaction/distinct_helpers_nested_nan_unequal_any.cu @@ -9,10 +9,11 @@ #include #include -#include #include #include +#include + namespace cudf::detail { template rmm::device_uvector reduce_by_row_keep_any( @@ -22,7 +23,7 @@ template rmm::device_uvector reduce_by_row_keep_any( cudf::detail::row::equality::physical_equality_comparator>, distinct_precomputed_hash>& set, size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace cudf::detail diff --git a/cpp/src/stream_compaction/distinct_helpers_nested_nan_unequal_ordered.cu b/cpp/src/stream_compaction/distinct_helpers_nested_nan_unequal_ordered.cu index 0a1b98929e81..a2451bfa4e24 100644 --- a/cpp/src/stream_compaction/distinct_helpers_nested_nan_unequal_ordered.cu +++ b/cpp/src/stream_compaction/distinct_helpers_nested_nan_unequal_ordered.cu @@ -10,10 +10,11 @@ #include #include -#include #include #include +#include + namespace cudf::detail { template rmm::device_uvector reduce_by_row_keep_first_last_none( @@ -23,7 +24,7 @@ template rmm::device_uvector reduce_by_row_keep_first_last_none( cudf::detail::row::equality::physical_equality_comparator>>& set, size_type num_rows, duplicate_keep_option keep, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace cudf::detail diff --git a/cpp/src/stream_compaction/drop_nans.cu b/cpp/src/stream_compaction/drop_nans.cu index dd01305f19ac..cb94e2577df1 100644 --- a/cpp/src/stream_compaction/drop_nans.cu +++ b/cpp/src/stream_compaction/drop_nans.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -14,8 +14,7 @@ #include #include -#include - +#include #include #include @@ -78,7 +77,7 @@ namespace detail { std::unique_ptr
drop_nans(table_view const& input, std::vector const& keys, cudf::size_type keep_threshold, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto keys_view = input.select(keys); @@ -106,7 +105,7 @@ std::unique_ptr
drop_nans(table_view const& input, std::unique_ptr
drop_nans(table_view const& input, std::vector const& keys, cudf::size_type keep_threshold, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -117,7 +116,7 @@ std::unique_ptr
drop_nans(table_view const& input, */ std::unique_ptr
drop_nans(table_view const& input, std::vector const& keys, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/stream_compaction/drop_nulls.cu b/cpp/src/stream_compaction/drop_nulls.cu index 24f45809e092..83c0724370a0 100644 --- a/cpp/src/stream_compaction/drop_nulls.cu +++ b/cpp/src/stream_compaction/drop_nulls.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -13,8 +13,7 @@ #include #include -#include - +#include #include #include @@ -57,7 +56,7 @@ namespace detail { std::unique_ptr
drop_nulls(table_view const& input, std::vector const& keys, cudf::size_type keep_threshold, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto keys_view = input.select(keys); @@ -79,7 +78,7 @@ std::unique_ptr
drop_nulls(table_view const& input, std::unique_ptr
drop_nulls(table_view const& input, std::vector const& keys, cudf::size_type keep_threshold, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -90,7 +89,7 @@ std::unique_ptr
drop_nulls(table_view const& input, */ std::unique_ptr
drop_nulls(table_view const& input, std::vector const& keys, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/stream_compaction/filter/filter.cu b/cpp/src/stream_compaction/filter/filter.cu index c90221d8331d..67cac939a68f 100644 --- a/cpp/src/stream_compaction/filter/filter.cu +++ b/cpp/src/stream_compaction/filter/filter.cu @@ -12,7 +12,7 @@ #include #include -#include +#include #include @@ -29,7 +29,7 @@ std::unique_ptr
filter(std::string const& predicate_udf, std::span predicate_inputs, table_view const& filter_table, output_nullability predicate_nullability, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(filter_table.num_columns() > 0, @@ -73,7 +73,7 @@ std::unique_ptr
filter(std::string const& predicate_udf, std::unique_ptr
filter(table_view const& predicate_table, ast::expression const& predicate_expr, table_view const& filter_table, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto args = cudf::detail::row_ir::ast_converter::filter(cudf::detail::row_ir::target::CUDA, @@ -103,7 +103,7 @@ std::vector> filter_extended( std::optional user_data, null_aware is_null_aware, output_nullability predicate_nullability, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -126,7 +126,7 @@ std::vector> filter(std::vector const& pred std::optional user_data, null_aware is_null_aware, output_nullability predicate_nullability, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // legacy behavior was to detect which column were scalars based on their sizes diff --git a/cpp/src/stream_compaction/stable_distinct.cu b/cpp/src/stream_compaction/stable_distinct.cu index 869d46a9c8be..11269e0e7ecf 100644 --- a/cpp/src/stream_compaction/stable_distinct.cu +++ b/cpp/src/stream_compaction/stable_distinct.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -23,7 +23,7 @@ std::unique_ptr
stable_distinct(table_view const& input, duplicate_keep_option keep, null_equality nulls_equal, nan_equality nans_equal, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.num_rows() == 0 or input.num_columns() == 0 or keys.empty()) { @@ -71,7 +71,7 @@ std::unique_ptr
stable_distinct(table_view const& input, duplicate_keep_option keep, null_equality nulls_equal, nan_equality nans_equal, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/stream_compaction/stream_compaction_common.cuh b/cpp/src/stream_compaction/stream_compaction_common.cuh index b0f0364850f5..70230ec39323 100644 --- a/cpp/src/stream_compaction/stream_compaction_common.cuh +++ b/cpp/src/stream_compaction/stream_compaction_common.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -7,10 +7,9 @@ #include #include -#include - #include #include +#include namespace cudf { namespace detail { @@ -68,7 +67,7 @@ OutputIterator unique_copy(InputIterator first, OutputIterator output, BinaryPredicate comp, duplicate_keep_option const keep, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { size_type const last_index = cuda::std::distance(first, last) - 1; return cudf::detail::copy_if( diff --git a/cpp/src/stream_compaction/unique.cu b/cpp/src/stream_compaction/unique.cu index e0da160b1d66..6b1d90b5e890 100644 --- a/cpp/src/stream_compaction/unique.cu +++ b/cpp/src/stream_compaction/unique.cu @@ -24,12 +24,12 @@ #include #include -#include #include #include #include #include +#include #include #include @@ -41,7 +41,7 @@ std::unique_ptr
unique(table_view const& input, std::vector const& keys, duplicate_keep_option keep, null_equality nulls_equal, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // If keep is KEEP_ANY, just alias it to KEEP_FIRST. @@ -112,7 +112,7 @@ std::unique_ptr
unique(table_view const& input, std::vector const& keys, duplicate_keep_option const keep, null_equality nulls_equal, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/table/table.cpp b/cpp/src/table/table.cpp index 1cfa1f41efdf..dd8f31113abb 100644 --- a/cpp/src/table/table.cpp +++ b/cpp/src/table/table.cpp @@ -8,14 +8,14 @@ #include #include -#include +#include #include namespace cudf { // Copy the columns from another table -table::table(table const& other, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) +table::table(table const& other, cuda::stream_ref stream, rmm::device_async_resource_ref mr) : _num_rows{other.num_rows()} { CUDF_FUNC_RANGE(); @@ -55,7 +55,7 @@ table::table(std::vector>&& columns, size_type num_rows) } // Copy the contents of a `table_view` -table::table(table_view view, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) +table::table(table_view view, cuda::stream_ref stream, rmm::device_async_resource_ref mr) : _num_rows{view.num_rows()} { CUDF_FUNC_RANGE(); diff --git a/cpp/src/table/table_device_view.cu b/cpp/src/table/table_device_view.cu index d179ef0b11c2..59f68f2d5d44 100644 --- a/cpp/src/table/table_device_view.cu +++ b/cpp/src/table/table_device_view.cu @@ -11,9 +11,10 @@ #include #include -#include #include +#include + #include #include @@ -42,7 +43,7 @@ template class table_device_view_base std::pair, ColumnDeviceView*> create_column_device_views( - HostTableView source_view, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + HostTableView source_view, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // First calculate the size of memory needed to hold the // table's ColumnDeviceViews. This is done by calling extent() @@ -82,7 +83,7 @@ std::pair, ColumnDeviceView*> create_column_ template std::pair, column_device_view*> create_column_device_views>( host_span source_view, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); table_device_view::table_device_view(table_view source_view, column_device_view* columns) @@ -92,7 +93,7 @@ table_device_view::table_device_view(table_view source_view, column_device_view* std::unique_ptr> table_device_view::create(table_view source_view, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto [descendant_storage, columns] = @@ -115,7 +116,7 @@ mutable_table_device_view::mutable_table_device_view(mutable_table_view source_v std::unique_ptr> mutable_table_device_view::create(mutable_table_view source_view, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto [descendant_storage, columns] = diff --git a/cpp/src/table/table_equal.cu b/cpp/src/table/table_equal.cu index 8033d2d408a5..74ef68bda46d 100644 --- a/cpp/src/table/table_equal.cu +++ b/cpp/src/table/table_equal.cu @@ -13,13 +13,13 @@ #include #include -#include #include #include #include #include #include +#include namespace cudf { namespace detail { @@ -29,7 +29,7 @@ template [[nodiscard]] bool tables_equal(table_view const& left, table_view const& right, null_equality nulls_equal, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const temp_mr = cudf::get_current_device_resource_ref(); auto const comparator = detail::row::equality::two_table_comparator{left, right, stream, temp_mr}; @@ -43,7 +43,7 @@ template [rows_equal] __device__(size_type i) -> bool { return rows_equal(detail::row::lhs_index_type{i}, detail::row::rhs_index_type{i}); }, - stream.value())); + stream.get())); return cudf::detail::reduce( eq_rows.begin(), eq_rows.end(), true, cuda::std::logical_and{}, stream); } @@ -53,7 +53,7 @@ template [[nodiscard]] bool tables_equal(table_view const& left, table_view const& right, null_equality nulls_equal, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (left.num_rows() != right.num_rows() || left.num_columns() != right.num_columns() || !have_same_types(left, right)) { @@ -71,7 +71,7 @@ template bool tables_equal(table_view const& left, table_view const& right, null_equality nulls_equal, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); return detail::tables_equal(left, right, nulls_equal, stream); diff --git a/cpp/src/transpose/transpose.cu b/cpp/src/transpose/transpose.cu index 42cf01c244ed..e842de457f93 100644 --- a/cpp/src/transpose/transpose.cu +++ b/cpp/src/transpose/transpose.cu @@ -15,14 +15,14 @@ #include #include -#include - #include +#include +#include namespace cudf { namespace detail { std::pair, table_view> transpose(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // If there are no rows in the input, return successfully @@ -49,7 +49,7 @@ std::pair, table_view> transpose(table_view const& input } // namespace detail std::pair, table_view> transpose(table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/utilities/cuda_memcpy.cu b/cpp/src/utilities/cuda_memcpy.cu index 38ed7231a062..045bcfe046fc 100644 --- a/cpp/src/utilities/cuda_memcpy.cu +++ b/cpp/src/utilities/cuda_memcpy.cu @@ -28,7 +28,7 @@ CUDF_KERNEL void copy_kernel(char const* __restrict__ src, char* __restrict__ ds if (idx < n) { dst[idx] = src[idx]; } } -void copy_pinned(void* dst, void const* src, std::size_t size, rmm::cuda_stream_view stream) +void copy_pinned(void* dst, void const* src, std::size_t size, cuda::stream_ref stream) { if (size == 0) return; @@ -37,32 +37,41 @@ void copy_pinned(void* dst, void const* src, std::size_t size, rmm::cuda_stream_ auto const grid_size = cudf::util::div_rounding_up_safe(size, block_size); // We are explicitly launching the kernel here instead of calling a thrust function because the // thrust function can potentially call cudaMemcpyAsync instead of using a kernel - copy_kernel<<>>( + copy_kernel<<>>( static_cast(src), static_cast(dst), size); } else { CUDF_CUDA_TRY(cudf::detail::memcpy_async(dst, src, size, stream)); } } -void copy_pageable(void* dst, void const* src, std::size_t size, rmm::cuda_stream_view stream) +void copy_pageable(void* dst, void const* src, std::size_t size, cuda::stream_ref stream) { if (size == 0) return; CUDF_CUDA_TRY(cudf::detail::memcpy_async(dst, src, size, stream)); } +#if CUDART_VERSION >= 13000 +bool is_default_stream(cuda::stream_ref stream) +{ + auto const cstream = stream.get(); + return cstream == cudaStreamDefault || cstream == cudaStreamLegacy || + cstream == cudaStreamPerThread; +} +#endif // CUDART_VERSION >= 13000 + }; // namespace cudaError_t memcpy_batch_async(void* const* dsts, void const* const* srcs, std::size_t const* sizes, std::size_t count, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // Uses cudaMemcpyBatchAsync for CUDA 13.0+ to avoid driver-side locking overhead. // cudaMemcpyBatchAsync does not support the default stream. #if CUDART_VERSION >= 13000 - if (!stream.is_default()) { + if (!is_default_stream(stream)) { constexpr std::size_t prefer_overlap_threshold = 128 * 1024; // Filter out invalid copies (nullptr dst/src or size==0); @@ -98,18 +107,18 @@ cudaError_t memcpy_batch_async(void* const* dsts, : cudaMemcpyFlagPreferOverlapWithCompute; cudaMemcpyAttributes attrs = {.srcAccessOrder = cudaMemcpySrcAccessOrderStream, .flags = flags}; std::size_t attrs_idx = 0; - return cudaMemcpyBatchAsync(dsts, srcs, sizes, count, &attrs, &attrs_idx, 1, stream.value()); + return cudaMemcpyBatchAsync(dsts, srcs, sizes, count, &attrs, &attrs_idx, 1, stream.get()); } #endif // CUDART_VERSION >= 13000 for (std::size_t i = 0; i < count; ++i) { cudaError_t status = - cudaMemcpyAsync(dsts[i], srcs[i], sizes[i], cudaMemcpyDefault, stream.value()); + cudaMemcpyAsync(dsts[i], srcs[i], sizes[i], cudaMemcpyDefault, stream.get()); if (status != cudaSuccess) { return status; } } return cudaSuccess; } -cudaError_t memcpy_async(void* dst, void const* src, size_t count, rmm::cuda_stream_view stream) +cudaError_t memcpy_async(void* dst, void const* src, size_t count, cuda::stream_ref stream) { if (count == 0) { return cudaSuccess; } @@ -119,7 +128,7 @@ cudaError_t memcpy_async(void* dst, void const* src, size_t count, rmm::cuda_str } void cuda_memcpy_async_impl( - void* dst, void const* src, size_t size, host_memory_kind kind, rmm::cuda_stream_view stream) + void* dst, void const* src, size_t size, host_memory_kind kind, cuda::stream_ref stream) { if (kind == host_memory_kind::PINNED) { copy_pinned(dst, src, size, stream); diff --git a/cpp/src/utilities/prefetch.cpp b/cpp/src/utilities/prefetch.cpp index 0f1e44a6a42e..3ec133f0d437 100644 --- a/cpp/src/utilities/prefetch.cpp +++ b/cpp/src/utilities/prefetch.cpp @@ -29,7 +29,7 @@ std::atomic_bool& debug() cudaError_t prefetch_noexcept(void const* ptr, std::size_t size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::cuda_device_id device_id) noexcept { if (!detail::enabled()) { return cudaSuccess; } @@ -53,9 +53,9 @@ cudaError_t prefetch_noexcept(void const* ptr, (device_id.value() == cudaCpuDeviceId) ? cudaMemLocationTypeHost : cudaMemLocationTypeDevice, {device_id.value()}}; constexpr int flags = 0; - auto result = cudaMemPrefetchAsync(ptr, size, location, flags, stream.value()); + auto result = cudaMemPrefetchAsync(ptr, size, location, flags, stream.get()); #else - auto result = cudaMemPrefetchAsync(ptr, size, device_id.value(), stream.value()); + auto result = cudaMemPrefetchAsync(ptr, size, device_id.value(), stream.get()); #endif // Need to flush the CUDA error so that the context is not corrupted. if (result == cudaErrorInvalidValue) { cudaGetLastError(); } @@ -64,7 +64,7 @@ cudaError_t prefetch_noexcept(void const* ptr, void prefetch(void const* ptr, std::size_t size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::cuda_device_id device_id) { auto result = prefetch_noexcept(ptr, size, stream, device_id); diff --git a/cpp/src/utilities/roaring_bitmap.cu b/cpp/src/utilities/roaring_bitmap.cu index 436a9b9ba46c..aea22e20e0c2 100644 --- a/cpp/src/utilities/roaring_bitmap.cu +++ b/cpp/src/utilities/roaring_bitmap.cu @@ -53,7 +53,7 @@ class roaring_bitmap::impl { impl& operator=(impl&&) = default; template - void materialize(rmm::cuda_stream_view stream) + void materialize(cuda::stream_ref stream) { auto const bytes = _serialized_bitmap_data.data(); @@ -89,7 +89,7 @@ class roaring_bitmap::impl { } template - void contains_async(InputIt first, InputIt last, OutputIt output, rmm::cuda_stream_view stream) + void contains_async(InputIt first, InputIt last, OutputIt output, cuda::stream_ref stream) { if (first == last) { return; } @@ -140,7 +140,7 @@ roaring_bitmap::roaring_bitmap(roaring_bitmap&&) noexcept = default; roaring_bitmap& roaring_bitmap::operator=(roaring_bitmap&&) noexcept = default; -void roaring_bitmap::materialize(rmm::cuda_stream_view stream) const +void roaring_bitmap::materialize(cuda::stream_ref stream) const { dispatch_roaring_bitmap_type( _type, [&]() { _impl->materialize(stream); }); @@ -167,9 +167,7 @@ cuda::std::size_t roaring_bitmap::size_bytes() const } std::unique_ptr roaring_bitmap::contains_async( - cudf::column_view const& keys, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) const + cudf::column_view const& keys, cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { auto result = cudf::make_fixed_width_column( cudf::data_type{cudf::type_id::BOOL8}, keys.size(), cudf::mask_state::UNALLOCATED, stream, mr); @@ -179,7 +177,7 @@ std::unique_ptr roaring_bitmap::contains_async( void roaring_bitmap::contains_async(cudf::column_view const& keys, cudf::mutable_column_view const& output, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { CUDF_EXPECTS(output.type().id() == cudf::type_id::BOOL8, "Output column must be BOOL8",