diff --git a/cpp/include/cudf/column/column.hpp b/cpp/include/cudf/column/column.hpp index 6f9f9b0d67f9..d02d821d83cd 100644 --- a/cpp/include/cudf/column/column.hpp +++ b/cpp/include/cudf/column/column.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 @@ -10,10 +10,11 @@ #include #include -#include #include #include +#include + #include #include #include @@ -52,7 +53,7 @@ class column { * @param mr Device memory resource to use for all device memory allocations */ column(column 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()); /** @@ -130,7 +131,7 @@ class column { * @param mr Device memory resource to use for all device memory allocations */ explicit column(column_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()); /** @@ -181,7 +182,7 @@ class column { */ void set_null_mask(rmm::device_buffer const& new_null_mask, size_type new_null_count, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Updates the count of null elements. diff --git a/cpp/include/cudf/column/column_device_view.cuh b/cpp/include/cudf/column/column_device_view.cuh index 756edd3a5f28..a7add0c42c26 100644 --- a/cpp/include/cudf/column/column_device_view.cuh +++ b/cpp/include/cudf/column/column_device_view.cuh @@ -15,11 +15,11 @@ #include #include -#include #include #include #include +#include #include #include @@ -506,7 +506,7 @@ class alignas(16) column_device_view : public column_device_view_core { */ static std::unique_ptr> create( column_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()); /** @@ -657,7 +657,7 @@ class alignas(16) mutable_column_device_view : public mutable_column_device_view static std::unique_ptr> create(mutable_column_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()); /** diff --git a/cpp/include/cudf/column/column_factories.hpp b/cpp/include/cudf/column/column_factories.hpp index 3423c89d0cee..b8b4b3fa5781 100644 --- a/cpp/include/cudf/column/column_factories.hpp +++ b/cpp/include/cudf/column/column_factories.hpp @@ -14,6 +14,7 @@ #include #include +#include /** * @file @@ -68,7 +69,7 @@ std::unique_ptr make_numeric_column( data_type type, size_type size, mask_state state = mask_state::UNALLOCATED, - 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 +95,7 @@ std::unique_ptr make_numeric_column( size_type size, B&& null_mask, size_type null_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()) { CUDF_EXPECTS(is_numeric(type), "Invalid, non-numeric type."); @@ -126,7 +127,7 @@ std::unique_ptr make_fixed_point_column( data_type type, size_type size, mask_state state = mask_state::UNALLOCATED, - 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()); /** @@ -151,7 +152,7 @@ std::unique_ptr make_fixed_point_column( size_type size, B&& null_mask, size_type null_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()) { CUDF_EXPECTS(is_fixed_point(type), "Invalid, non-fixed_point type."); @@ -184,7 +185,7 @@ std::unique_ptr make_timestamp_column( data_type type, size_type size, mask_state state = mask_state::UNALLOCATED, - 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()); /** @@ -210,7 +211,7 @@ std::unique_ptr make_timestamp_column( size_type size, B&& null_mask, size_type null_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()) { CUDF_EXPECTS(is_timestamp(type), "Invalid, non-timestamp type."); @@ -243,7 +244,7 @@ std::unique_ptr make_duration_column( data_type type, size_type size, mask_state state = mask_state::UNALLOCATED, - 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()); /** @@ -269,7 +270,7 @@ std::unique_ptr make_duration_column( size_type size, B&& null_mask, size_type null_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()) { CUDF_EXPECTS(is_duration(type), "Invalid, non-duration type."); @@ -302,7 +303,7 @@ std::unique_ptr make_fixed_width_column( data_type type, size_type size, mask_state state = mask_state::UNALLOCATED, - 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()); /** @@ -328,7 +329,7 @@ std::unique_ptr make_fixed_width_column( size_type size, B&& null_mask, size_type null_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()) { CUDF_EXPECTS(is_fixed_width(type), "Invalid, non-fixed-width type."); @@ -593,7 +594,7 @@ std::unique_ptr create_structs_hierarchy( std::unique_ptr make_column_from_scalar( scalar const& s, size_type size, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -613,7 +614,7 @@ std::unique_ptr make_column_from_scalar( std::unique_ptr make_dictionary_from_scalar( scalar const& s, size_type size, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @} */ // end of group diff --git a/cpp/include/cudf/column/column_stream.hpp b/cpp/include/cudf/column/column_stream.hpp index ac2e1f622f62..db1835adc8dd 100644 --- a/cpp/include/cudf/column/column_stream.hpp +++ b/cpp/include/cudf/column/column_stream.hpp @@ -6,7 +6,7 @@ #include -#include +#include #include @@ -40,7 +40,7 @@ namespace CUDF_EXPORT cudf { * @param stream Stream used for future asynchronous deallocation of the buffers * @return Column with equivalent contents and rebinding applied */ -[[nodiscard]] std::unique_ptr rebind_stream(column&& col, rmm::cuda_stream_view stream); +[[nodiscard]] std::unique_ptr rebind_stream(column&& col, cuda::stream_ref stream); /** @} */ // end of group } // namespace CUDF_EXPORT cudf diff --git a/cpp/include/cudf/column/column_view.hpp b/cpp/include/cudf/column/column_view.hpp index 24a020f257aa..f8dd55c76e02 100644 --- a/cpp/include/cudf/column/column_view.hpp +++ b/cpp/include/cudf/column/column_view.hpp @@ -170,10 +170,9 @@ class column_view_base { * @param[in] stream CUDA stream used for device memory operations and kernel launches * @return The count of null elements in the given range */ - [[nodiscard]] size_type null_count( - size_type begin, - size_type end, - rmm::cuda_stream_view stream = cudf::get_default_stream()) const; + [[nodiscard]] size_type null_count(size_type begin, + size_type end, + cuda::stream_ref stream = cudf::get_default_stream()) const; /** * @brief Indicates if the column contains null elements, @@ -199,7 +198,7 @@ class column_view_base { */ [[nodiscard]] bool has_nulls(size_type begin, size_type end, - rmm::cuda_stream_view stream = cudf::get_default_stream()) const + cuda::stream_ref stream = cudf::get_default_stream()) const { return null_count(begin, end, stream) > 0; } diff --git a/cpp/include/cudf/concatenate.hpp b/cpp/include/cudf/concatenate.hpp index a78b62cca2f6..8f6459d74247 100644 --- a/cpp/include/cudf/concatenate.hpp +++ b/cpp/include/cudf/concatenate.hpp @@ -11,6 +11,8 @@ #include #include +#include + #include #include @@ -39,7 +41,7 @@ namespace CUDF_EXPORT cudf { */ rmm::device_buffer concatenate_masks( std::span views, - 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()); /** @@ -56,7 +58,7 @@ rmm::device_buffer concatenate_masks( */ std::unique_ptr concatenate( std::span columns_to_concat, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -85,7 +87,7 @@ std::unique_ptr concatenate( */ std::unique_ptr concatenate( std::span tables_to_concat, - 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/contiguous_split.hpp b/cpp/include/cudf/contiguous_split.hpp index e98cac984d3f..673a4bb9feee 100644 --- a/cpp/include/cudf/contiguous_split.hpp +++ b/cpp/include/cudf/contiguous_split.hpp @@ -70,7 +70,7 @@ namespace CUDF_EXPORT cudf { std::vector contiguous_split( cudf::table_view const& input, std::vector const& splits, - 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 { @@ -155,7 +155,7 @@ class chunked_pack { explicit chunked_pack( cudf::table_view const& input, std::size_t user_buffer_size, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref temp_mr = cudf::get_current_device_resource_ref()); /** @@ -222,7 +222,7 @@ class chunked_pack { [[nodiscard]] static std::unique_ptr create( cudf::table_view const& input, std::size_t user_buffer_size, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref temp_mr = cudf::get_current_device_resource_ref()); private: @@ -244,7 +244,7 @@ class chunked_pack { * and device memory respectively */ packed_columns pack(cudf::table_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -261,7 +261,7 @@ packed_columns pack(cudf::table_view const& input, */ std::size_t packed_size( cudf::table_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref temp_mr = cudf::get_current_device_resource_ref()); /** diff --git a/cpp/include/cudf/copying.hpp b/cpp/include/cudf/copying.hpp index b744d65cb25c..f1ab6e30f268 100644 --- a/cpp/include/cudf/copying.hpp +++ b/cpp/include/cudf/copying.hpp @@ -82,7 +82,7 @@ std::unique_ptr
gather( table_view const& source_table, column_view const& gather_map, out_of_bounds_policy bounds_policy = out_of_bounds_policy::DONT_CHECK, - 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()); /** @@ -120,7 +120,7 @@ std::unique_ptr
gather( column_view const& gather_map, out_of_bounds_policy bounds_policy, negative_index_policy neg_indices, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -140,7 +140,7 @@ std::unique_ptr
gather( */ std::unique_ptr
reverse( table_view const& source_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()); /** @@ -160,7 +160,7 @@ std::unique_ptr
reverse( */ std::unique_ptr reverse( column_view const& source_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()); /** @@ -209,7 +209,7 @@ std::unique_ptr
scatter( table_view const& source, column_view const& scatter_map, table_view const& target, - 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()); /** @@ -252,7 +252,7 @@ std::unique_ptr
scatter( std::vector> const& source, column_view const& indices, table_view const& target, - 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()); /** @@ -300,7 +300,7 @@ std::unique_ptr empty_like(scalar const& input); std::unique_ptr allocate_like( column_view const& input, mask_allocation_policy mask_alloc = mask_allocation_policy::RETAIN, - 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()); /** @@ -323,7 +323,7 @@ std::unique_ptr allocate_like( column_view const& input, size_type size, mask_allocation_policy mask_alloc = mask_allocation_policy::RETAIN, - 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 @@ void copy_range_in_place(column_view const& source, size_type source_begin, size_type source_end, size_type target_begin, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Copies a range of elements out-of-place from one column to another. @@ -415,7 +415,7 @@ std::unique_ptr copy_range( size_type source_begin, size_type source_end, size_type target_begin, - 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()); /** @@ -459,7 +459,7 @@ std::unique_ptr shift( column_view const& input, size_type offset, scalar const& fill_value, - 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()); /** @@ -495,14 +495,14 @@ std::unique_ptr shift( */ std::vector slice(column_view const& input, std::span indices, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @ingroup copy_slice - * @copydoc cudf::slice(column_view const&, std::span, rmm::cuda_stream_view) + * @copydoc cudf::slice(column_view const&, std::span, cuda::stream_ref) */ std::vector slice(column_view const& input, std::initializer_list indices, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Slices a `table_view` into a set of `table_view`s according to a set of indices. @@ -539,14 +539,14 @@ std::vector slice(column_view const& input, */ std::vector slice(table_view const& input, std::span indices, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @ingroup copy_slice - * @copydoc cudf::slice(table_view const&, std::span, rmm::cuda_stream_view stream) + * @copydoc cudf::slice(table_view const&, std::span, cuda::stream_ref stream) */ std::vector slice(table_view const& input, std::initializer_list indices, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Splits a `column_view` into a set of `column_view`s according to a set of indices @@ -583,14 +583,14 @@ std::vector slice(table_view const& input, */ std::vector split(column_view const& input, std::span splits, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @ingroup copy_split - * @copydoc cudf::split(column_view const&, std::span, rmm::cuda_stream_view) + * @copydoc cudf::split(column_view const&, std::span, cuda::stream_ref) */ std::vector split(column_view const& input, std::initializer_list splits, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Splits a `table_view` into a set of `table_view`s according to a set of indices @@ -629,14 +629,14 @@ std::vector split(column_view const& input, */ std::vector split(table_view const& input, std::span splits, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @ingroup copy_split - * @copydoc cudf::split(table_view const&, std::span, rmm::cuda_stream_view) + * @copydoc cudf::split(table_view const&, std::span, cuda::stream_ref) */ std::vector split(table_view const& input, std::initializer_list splits, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Returns a new column, where each element is selected from either @p lhs or @@ -662,7 +662,7 @@ std::unique_ptr copy_if_else( column_view const& lhs, column_view const& rhs, 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()); /** @@ -688,7 +688,7 @@ std::unique_ptr copy_if_else( scalar const& lhs, column_view const& rhs, 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()); /** @@ -714,7 +714,7 @@ std::unique_ptr copy_if_else( column_view const& lhs, scalar const& rhs, 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()); /** @@ -738,7 +738,7 @@ std::unique_ptr copy_if_else( scalar const& lhs, scalar const& rhs, 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()); /** @@ -782,7 +782,7 @@ std::unique_ptr
boolean_mask_scatter( table_view const& input, table_view const& target, 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()); /** @@ -821,7 +821,7 @@ std::unique_ptr
boolean_mask_scatter( std::vector> const& input, table_view const& target, 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()); /** @@ -841,7 +841,7 @@ std::unique_ptr
boolean_mask_scatter( std::unique_ptr get_element( column_view const& input, size_type index, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -885,7 +885,7 @@ std::unique_ptr
sample( size_type const n, sample_with_replacement replacement = sample_with_replacement::FALSE, int64_t const seed = 0, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -906,7 +906,7 @@ std::unique_ptr
sample( * @return false If neither the column or its descendants have non-empty null rows */ bool has_nonempty_nulls(column_view const& input, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Approximates if a column or its descendants *may* have non-empty null elements @@ -1002,7 +1002,7 @@ bool may_have_nonempty_nulls(column_view const& input); */ std::unique_ptr purge_nonempty_nulls( 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()); /** @} */ diff --git a/cpp/include/cudf/detail/concatenate.hpp b/cpp/include/cudf/detail/concatenate.hpp index 5056a05732a3..b67f556b911d 100644 --- a/cpp/include/cudf/detail/concatenate.hpp +++ b/cpp/include/cudf/detail/concatenate.hpp @@ -11,7 +11,7 @@ #include #include -#include +#include #include #include @@ -21,18 +21,18 @@ namespace cudf { namespace detail { /** * @copydoc cudf::concatenate(std::span,rmm::cuda_stream_view,rmm::device_async_resource_ref) + * const>,cuda::stream_ref,rmm::device_async_resource_ref) */ std::unique_ptr concatenate(std::span columns_to_concat, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @copydoc cudf::concatenate(std::span,rmm::cuda_stream_view,rmm::device_async_resource_ref) + * const>,cuda::stream_ref,rmm::device_async_resource_ref) */ std::unique_ptr
concatenate(std::span tables_to_concat, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/detail/concatenate_masks.hpp b/cpp/include/cudf/detail/concatenate_masks.hpp index 8d4d1f272a08..449791bee14c 100644 --- a/cpp/include/cudf/detail/concatenate_masks.hpp +++ b/cpp/include/cudf/detail/concatenate_masks.hpp @@ -9,9 +9,10 @@ #include #include -#include #include +#include + #include namespace cudf { @@ -33,7 +34,7 @@ size_type concatenate_masks(device_span d_views, device_span d_offsets, bitmask_type* dest_mask, size_type output_size, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @brief Concatenates `views[i]`'s bitmask from the bits @@ -47,7 +48,7 @@ size_type concatenate_masks(device_span d_views, */ size_type concatenate_masks(host_span views, bitmask_type* dest_mask, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @copydoc cudf::concatenate_masks(std::span, rmm::device_async_resource_ref) @@ -55,7 +56,7 @@ size_type concatenate_masks(host_span views, * @param stream CUDA stream used for device memory operations and kernel launches. */ rmm::device_buffer concatenate_masks(std::span views, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/detail/contiguous_split.hpp b/cpp/include/cudf/detail/contiguous_split.hpp index 22d01417f048..73de92bce2e9 100644 --- a/cpp/include/cudf/detail/contiguous_split.hpp +++ b/cpp/include/cudf/detail/contiguous_split.hpp @@ -10,7 +10,7 @@ #include #include -#include +#include #include #include @@ -23,14 +23,14 @@ namespace detail { **/ std::vector contiguous_split(cudf::table_view const& input, std::vector const& splits, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @copydoc cudf::pack **/ packed_columns pack(cudf::table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); // opaque implementation of `metadata_builder` since it needs to use diff --git a/cpp/include/cudf/detail/copy.hpp b/cpp/include/cudf/detail/copy.hpp index 5037edb16e8e..0afd97f4b44e 100644 --- a/cpp/include/cudf/detail/copy.hpp +++ b/cpp/include/cudf/detail/copy.hpp @@ -13,7 +13,7 @@ #include #include -#include +#include #include #include @@ -39,10 +39,7 @@ namespace detail { * @return ColumnView View of the elements `[begin,end)` from `input`. */ template -ColumnView slice(ColumnView const& input, - size_type begin, - size_type end, - rmm::cuda_stream_view stream); +ColumnView slice(ColumnView const& input, size_type begin, size_type end, cuda::stream_ref stream); /** * @copydoc cudf::slice(column_view const&, std::span) @@ -51,7 +48,7 @@ ColumnView slice(ColumnView const& input, */ std::vector slice(column_view const& input, std::span indices, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @copydoc cudf::slice(column_view const&, std::initializer_list) * @@ -59,7 +56,7 @@ std::vector slice(column_view const& input, */ std::vector slice(column_view const& input, std::initializer_list indices, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @copydoc cudf::slice(table_view const&, std::span) @@ -68,7 +65,7 @@ std::vector slice(column_view const& input, */ std::vector slice(table_view const& input, std::span indices, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @copydoc cudf::slice(table_view const&, std::initializer_list) * @@ -76,7 +73,7 @@ std::vector slice(table_view const& input, */ std::vector slice(table_view const& input, std::initializer_list indices, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @copydoc cudf::split(column_view const&, std::span) @@ -85,7 +82,7 @@ std::vector slice(table_view const& input, */ std::vector split(column_view const& input, std::span splits, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @copydoc cudf::split(column_view const&, std::initializer_list) * @@ -93,7 +90,7 @@ std::vector split(column_view const& input, */ std::vector split(column_view const& input, std::initializer_list splits, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @copydoc cudf::split(table_view const&, std::span) @@ -102,7 +99,7 @@ std::vector split(column_view const& input, */ std::vector split(table_view const& input, std::span splits, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @copydoc cudf::split(table_view const&, std::initializer_list) * @@ -110,7 +107,7 @@ std::vector split(table_view const& input, */ std::vector split(table_view const& input, std::initializer_list splits, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @copydoc cudf::shift(column_view const&,size_type,scalar const&, @@ -121,7 +118,7 @@ std::vector split(table_view const& input, std::unique_ptr shift(column_view const& input, size_type offset, scalar const& fill_value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -161,7 +158,7 @@ std::unique_ptr segmented_shift(column_view const& segmented_values, device_span segment_offsets, size_type offset, scalar const& fill_value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -173,7 +170,7 @@ std::unique_ptr segmented_shift(column_view const& segmented_values, std::unique_ptr allocate_like(column_view const& input, size_type size, mask_allocation_policy mask_alloc, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -185,7 +182,7 @@ std::unique_ptr allocate_like(column_view const& input, std::unique_ptr copy_if_else(column_view const& lhs, column_view const& rhs, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -197,7 +194,7 @@ std::unique_ptr copy_if_else(column_view const& lhs, std::unique_ptr copy_if_else(scalar const& lhs, column_view const& rhs, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -209,7 +206,7 @@ std::unique_ptr copy_if_else(scalar const& lhs, std::unique_ptr copy_if_else(column_view const& lhs, scalar const& rhs, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -221,7 +218,7 @@ std::unique_ptr copy_if_else(column_view const& lhs, std::unique_ptr copy_if_else(scalar const& lhs, scalar const& rhs, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -233,7 +230,7 @@ std::unique_ptr
sample(table_view const& input, size_type const n, sample_with_replacement replacement, int64_t const seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -243,7 +240,7 @@ std::unique_ptr
sample(table_view const& input, */ std::unique_ptr get_element(column_view const& input, size_type index, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -251,14 +248,14 @@ std::unique_ptr get_element(column_view const& input, * * @param stream CUDA stream used for device memory operations and kernel launches. */ -bool has_nonempty_nulls(column_view const& input, rmm::cuda_stream_view stream); +bool has_nonempty_nulls(column_view const& input, cuda::stream_ref stream); /** * @copydoc cudf::may_have_nonempty_nulls * * @param stream CUDA stream used for device memory operations and kernel launches. */ -bool may_have_nonempty_nulls(column_view const& input, rmm::cuda_stream_view stream); +bool may_have_nonempty_nulls(column_view const& input, cuda::stream_ref stream); /** * @copydoc cudf::purge_nonempty_nulls @@ -266,7 +263,7 @@ bool may_have_nonempty_nulls(column_view const& input, rmm::cuda_stream_view str * @param stream CUDA stream used for device memory operations and kernel launches. */ std::unique_ptr purge_nonempty_nulls(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/copy_range.cuh b/cpp/include/cudf/detail/copy_range.cuh index 0e958749bfba..dcb934dd645e 100644 --- a/cpp/include/cudf/detail/copy_range.cuh +++ b/cpp/include/cudf/detail/copy_range.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 */ #pragma once @@ -18,9 +18,8 @@ #include #include -#include - #include +#include #include #include @@ -122,7 +121,7 @@ void copy_range(SourceValueIterator source_value_begin, mutable_column_view& target, size_type target_begin, size_type target_end, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS((target_begin <= target_end) && (target_begin >= 0) && (target_begin < target.size()) && (target_end <= target.size()), @@ -146,7 +145,7 @@ void copy_range(SourceValueIterator source_value_begin, auto kernel = copy_range_kernel; - kernel<<>>( + kernel<<>>( source_value_begin, source_validity_begin, *mutable_column_device_view::create(target, stream), @@ -158,7 +157,7 @@ void copy_range(SourceValueIterator source_value_begin, } else { auto kernel = copy_range_kernel; - kernel<<>>( + kernel<<>>( source_value_begin, source_validity_begin, *mutable_column_device_view::create(target, stream), @@ -167,7 +166,7 @@ void copy_range(SourceValueIterator source_value_begin, nullptr); } - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); } /** @@ -179,7 +178,7 @@ void copy_range_in_place(column_view const& source, size_type source_begin, size_type source_end, size_type target_begin, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @copydoc cudf::copy_range @@ -191,7 +190,7 @@ std::unique_ptr copy_range(column_view const& source, size_type source_begin, size_type source_end, size_type target_begin, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace cudf::detail diff --git a/cpp/include/cudf/detail/fill.hpp b/cpp/include/cudf/detail/fill.hpp index 7efe6ccbbd05..4931b2f69363 100644 --- a/cpp/include/cudf/detail/fill.hpp +++ b/cpp/include/cudf/detail/fill.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 */ @@ -10,7 +10,7 @@ #include #include -#include +#include #include @@ -24,7 +24,7 @@ void fill_in_place(mutable_column_view& destination, size_type begin, size_type end, scalar const& value, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @copydoc cudf::fill @@ -33,7 +33,7 @@ std::unique_ptr fill(column_view const& input, size_type begin, size_type end, scalar const& value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/detail/gather.hpp b/cpp/include/cudf/detail/gather.hpp index 747185cbbec1..9190b08e57f2 100644 --- a/cpp/include/cudf/detail/gather.hpp +++ b/cpp/include/cudf/detail/gather.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 @@ -13,7 +13,7 @@ #include #include -#include +#include #include @@ -22,19 +22,19 @@ namespace detail { /** * @copydoc cudf::gather(table_view const&,column_view const&,table_view - * const&,cudf::out_of_bounds_policy,cudf::negative_index_policy,rmm::cuda_stream_view, + * const&,cudf::out_of_bounds_policy,cudf::negative_index_policy,cuda::stream_ref, * rmm::device_async_resource_ref) */ std::unique_ptr
gather(table_view const& source_table, column_view const& gather_map, out_of_bounds_policy bounds_policy, negative_index_policy neg_indices, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @copydoc cudf::detail::gather(table_view const&,column_view const&,table_view - * const&,cudf::out_of_bounds_policy,cudf::negative_index_policy,rmm::cuda_stream_view, + * const&,cudf::out_of_bounds_policy,cudf::negative_index_policy,cuda::stream_ref, * rmm::device_async_resource_ref) * * @throws cudf::logic_error if `gather_map` span size is larger than max of `size_type`. @@ -43,7 +43,7 @@ std::unique_ptr
gather(table_view const& source_table, device_span const gather_map, out_of_bounds_policy bounds_policy, negative_index_policy neg_indices, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/detail/repeat.hpp b/cpp/include/cudf/detail/repeat.hpp index 967632570fde..9ef70c4f951e 100644 --- a/cpp/include/cudf/detail/repeat.hpp +++ b/cpp/include/cudf/detail/repeat.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 */ @@ -9,7 +9,7 @@ #include #include -#include +#include #include @@ -18,21 +18,21 @@ namespace detail { /** * @copydoc cudf::repeat(table_view const&, column_view const&, bool, - * rmm::cuda_stream_view, rmm::device_async_resource_ref) + * cuda::stream_ref, rmm::device_async_resource_ref) */ std::unique_ptr
repeat(table_view const& input_table, column_view const& count, bool check_count, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @copydoc cudf::repeat(table_view const&, size_type, - * rmm::cuda_stream_view, rmm::device_async_resource_ref) + * cuda::stream_ref, rmm::device_async_resource_ref) */ std::unique_ptr
repeat(table_view const& input_table, size_type count, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/detail/scatter.cuh b/cpp/include/cudf/detail/scatter.cuh index b2610b99f966..a836d7eb33eb 100644 --- a/cpp/include/cudf/detail/scatter.cuh +++ b/cpp/include/cudf/detail/scatter.cuh @@ -23,10 +23,10 @@ #include #include -#include #include #include +#include #include #include #include @@ -58,7 +58,7 @@ template auto scatter_to_gather(MapIterator scatter_map_begin, MapIterator scatter_map_end, size_type gather_rows, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { using MapValueType = cuda::std::iter_value_t; @@ -101,7 +101,7 @@ template auto scatter_to_gather_complement(MapIterator scatter_map_begin, MapIterator scatter_map_end, size_type gather_rows, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto gather_map = rmm::device_uvector(gather_rows, stream); thrust::sequence(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), @@ -137,7 +137,7 @@ struct column_scatterer_impl(target, stream, mr); @@ -162,7 +162,7 @@ struct column_scatterer_impl { MapIterator scatter_map_begin, MapIterator scatter_map_end, column_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { auto d_column = column_device_view::create(source, stream); @@ -179,7 +179,7 @@ struct column_scatterer_impl { MapIterator scatter_map_begin, MapIterator scatter_map_end, column_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { return cudf::lists::detail::scatter( @@ -194,7 +194,7 @@ struct column_scatterer_impl { MapIterator scatter_map_begin, MapIterator scatter_map_end, column_view const& target_in, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { if (target_in.is_empty()) // empty begets empty @@ -240,7 +240,7 @@ struct column_scatterer { MapIterator scatter_map_begin, MapIterator scatter_map_end, column_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { column_scatterer_impl scatterer{}; @@ -255,7 +255,7 @@ struct column_scatterer_impl { MapItRoot scatter_map_begin, MapItRoot scatter_map_end, column_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { CUDF_EXPECTS(source.num_children() == target.num_children(), @@ -368,7 +368,7 @@ std::unique_ptr
scatter(table_view const& source, MapIterator scatter_map_begin, MapIterator scatter_map_end, table_view const& target, - 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/scatter.hpp b/cpp/include/cudf/detail/scatter.hpp index bba6a767020e..f56424c298a6 100644 --- a/cpp/include/cudf/detail/scatter.hpp +++ b/cpp/include/cudf/detail/scatter.hpp @@ -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 */ #pragma once @@ -12,7 +12,7 @@ #include #include -#include +#include #include @@ -53,19 +53,19 @@ namespace detail { std::unique_ptr
scatter(table_view const& source, column_view const& scatter_map, table_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** * @copydoc cudf::detail::scatter(table_view const&,column_view const&,table_view - * const&,bool,rmm::cuda_stream_view,rmm::device_async_resource_ref) + * const&,bool,cuda::stream_ref,rmm::device_async_resource_ref) * * @throws cudf::logic_error if `scatter_map` span size is larger than max of `size_type`. */ std::unique_ptr
scatter(table_view const& source, device_span const scatter_map, table_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -100,7 +100,7 @@ std::unique_ptr
scatter(table_view const& source, std::unique_ptr
scatter(std::vector> const& source, column_view const& indices, table_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -114,7 +114,7 @@ std::unique_ptr
scatter(std::vector> std::unique_ptr
boolean_mask_scatter(table_view const& source, table_view const& target, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -130,7 +130,7 @@ std::unique_ptr
boolean_mask_scatter( std::vector> const& source, table_view const& target, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/detail/search.hpp b/cpp/include/cudf/detail/search.hpp index 2ee3d239deeb..8e9e80195a9e 100644 --- a/cpp/include/cudf/detail/search.hpp +++ b/cpp/include/cudf/detail/search.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 */ @@ -11,9 +11,10 @@ #include #include -#include #include +#include + namespace cudf { namespace detail { @@ -26,7 +27,7 @@ std::unique_ptr lower_bound(table_view const& haystack, table_view const& needles, 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); /** @@ -38,7 +39,7 @@ std::unique_ptr upper_bound(table_view const& haystack, table_view const& needles, 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); /** @@ -46,7 +47,7 @@ std::unique_ptr upper_bound(table_view const& haystack, * * @param stream CUDA stream used for device memory operations and kernel launches. */ -bool contains(column_view const& haystack, scalar const& needle, rmm::cuda_stream_view stream); +bool contains(column_view const& haystack, scalar const& needle, cuda::stream_ref stream); /** * @copydoc cudf::contains(column_view const&, column_view const&, rmm::device_async_resource_ref) @@ -55,7 +56,7 @@ bool contains(column_view const& haystack, scalar const& needle, rmm::cuda_strea */ std::unique_ptr contains(column_view const& haystack, column_view const& needles, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -88,7 +89,7 @@ rmm::device_uvector contains(table_view const& haystack, table_view const& needles, null_equality compare_nulls, nan_equality compare_nans, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/detail/sequence.hpp b/cpp/include/cudf/detail/sequence.hpp index 92c04eac722e..d67d5a83f4b9 100644 --- a/cpp/include/cudf/detail/sequence.hpp +++ b/cpp/include/cudf/detail/sequence.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,27 +10,27 @@ #include #include -#include +#include namespace cudf { namespace detail { /** - * @copydoc cudf::sequence(size_type, scalar const&, scalar const&,rmm::cuda_stream_view + * @copydoc cudf::sequence(size_type, scalar const&, scalar const&,cuda::stream_ref * stream,rmm::device_async_resource_ref) */ std::unique_ptr sequence(size_type size, scalar const& init, scalar const& step, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** - * @copydoc cudf::sequence(size_type, scalar const&, rmm::cuda_stream_view, + * @copydoc cudf::sequence(size_type, scalar const&, cuda::stream_ref, * rmm::device_async_resource_ref) */ std::unique_ptr sequence(size_type size, scalar const& init, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -42,7 +42,7 @@ std::unique_ptr sequence(size_type size, std::unique_ptr calendrical_month_sequence(size_type size, scalar const& init, size_type months, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/dictionary/detail/concatenate.hpp b/cpp/include/cudf/dictionary/detail/concatenate.hpp index 5dce586f205a..1238e3f6b1bc 100644 --- a/cpp/include/cudf/dictionary/detail/concatenate.hpp +++ b/cpp/include/cudf/dictionary/detail/concatenate.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 @@ -10,7 +10,7 @@ #include #include -#include +#include namespace cudf { namespace dictionary::detail { @@ -27,7 +27,7 @@ namespace dictionary::detail { * @return New column with concatenated results. */ std::unique_ptr concatenate(host_span columns, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace dictionary::detail diff --git a/cpp/include/cudf/dictionary/detail/encode.hpp b/cpp/include/cudf/dictionary/detail/encode.hpp index c5c0860871c9..e728acc1e389 100644 --- a/cpp/include/cudf/dictionary/detail/encode.hpp +++ b/cpp/include/cudf/dictionary/detail/encode.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 @@ -10,7 +10,7 @@ #include #include -#include +#include namespace cudf { namespace dictionary::detail { @@ -42,7 +42,7 @@ namespace dictionary::detail { */ std::unique_ptr encode(column_view const& column, data_type indices_type, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -61,7 +61,7 @@ std::unique_ptr encode(column_view const& column, * @return New column with type matching the dictionary_column's keys. */ std::unique_ptr decode(dictionary_column_view const& dictionary_column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** diff --git a/cpp/include/cudf/dictionary/detail/merge.hpp b/cpp/include/cudf/dictionary/detail/merge.hpp index 361000327aa0..de4f15a3a431 100644 --- a/cpp/include/cudf/dictionary/detail/merge.hpp +++ b/cpp/include/cudf/dictionary/detail/merge.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 @@ -9,7 +9,7 @@ #include #include -#include +#include namespace cudf { namespace dictionary::detail { @@ -32,7 +32,7 @@ namespace dictionary::detail { std::unique_ptr merge(dictionary_column_view const& lcol, dictionary_column_view const& rcol, cudf::detail::index_vector const& row_order, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace dictionary::detail diff --git a/cpp/include/cudf/dictionary/detail/replace.hpp b/cpp/include/cudf/dictionary/detail/replace.hpp index 3bab35ef8649..3cd91b0ce77b 100644 --- a/cpp/include/cudf/dictionary/detail/replace.hpp +++ b/cpp/include/cudf/dictionary/detail/replace.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 @@ -10,7 +10,7 @@ #include #include -#include +#include namespace cudf { namespace dictionary::detail { @@ -30,7 +30,7 @@ namespace dictionary::detail { */ std::unique_ptr replace_nulls(dictionary_column_view const& input, dictionary_column_view const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -47,7 +47,7 @@ std::unique_ptr replace_nulls(dictionary_column_view const& input, */ std::unique_ptr replace_nulls(dictionary_column_view const& input, scalar const& replacement, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace dictionary::detail diff --git a/cpp/include/cudf/dictionary/detail/search.hpp b/cpp/include/cudf/dictionary/detail/search.hpp index 8d908f1da30e..17483793ba44 100644 --- a/cpp/include/cudf/dictionary/detail/search.hpp +++ b/cpp/include/cudf/dictionary/detail/search.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 @@ -9,7 +9,7 @@ #include #include -#include +#include namespace cudf { namespace dictionary { @@ -23,7 +23,7 @@ namespace detail { */ std::unique_ptr get_index(dictionary_column_view const& dictionary, scalar const& key, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace detail diff --git a/cpp/include/cudf/dictionary/detail/update_keys.hpp b/cpp/include/cudf/dictionary/detail/update_keys.hpp index aa4de9355816..0fb608acfa02 100644 --- a/cpp/include/cudf/dictionary/detail/update_keys.hpp +++ b/cpp/include/cudf/dictionary/detail/update_keys.hpp @@ -11,7 +11,7 @@ #include #include -#include +#include #include @@ -25,7 +25,7 @@ namespace dictionary::detail { */ std::unique_ptr add_keys(dictionary_column_view const& dictionary_column, column_view const& new_keys, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -36,7 +36,7 @@ std::unique_ptr add_keys(dictionary_column_view const& dictionary_column */ std::unique_ptr remove_keys(dictionary_column_view const& dictionary_column, column_view const& keys_to_remove, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -46,7 +46,7 @@ std::unique_ptr remove_keys(dictionary_column_view const& dictionary_col * @param stream CUDA stream used for device memory operations and kernel launches. */ std::unique_ptr remove_unused_keys(dictionary_column_view const& dictionary_column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -57,7 +57,7 @@ std::unique_ptr remove_unused_keys(dictionary_column_view const& diction */ std::unique_ptr set_keys(dictionary_column_view const& dictionary_column, column_view const& keys, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -68,7 +68,7 @@ std::unique_ptr set_keys(dictionary_column_view const& dictionary_column */ std::vector> match_dictionaries( std::span input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); /** @@ -91,7 +91,7 @@ std::vector> match_dictionaries( * @return New dictionary columns and updated cudf::table_views. */ std::pair>, std::vector> match_dictionaries( - std::vector tables, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr); + std::vector tables, cuda::stream_ref stream, rmm::device_async_resource_ref mr); } // namespace dictionary::detail } // namespace cudf diff --git a/cpp/include/cudf/dictionary/dictionary_factories.hpp b/cpp/include/cudf/dictionary/dictionary_factories.hpp index 557d80958172..abfc29ba3e9f 100644 --- a/cpp/include/cudf/dictionary/dictionary_factories.hpp +++ b/cpp/include/cudf/dictionary/dictionary_factories.hpp @@ -9,7 +9,7 @@ #include #include -#include +#include /** * @file @@ -55,7 +55,7 @@ namespace CUDF_EXPORT cudf { std::unique_ptr make_dictionary_column( column_view const& keys_column, column_view const& indices_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()); /** @@ -103,7 +103,7 @@ std::unique_ptr make_dictionary_column(std::unique_ptr keys_colu std::unique_ptr make_dictionary_column( std::unique_ptr keys_column, std::unique_ptr indices_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()); /** @} */ // end of group diff --git a/cpp/include/cudf/dictionary/encode.hpp b/cpp/include/cudf/dictionary/encode.hpp index 555f9d85c376..d915927e7c0b 100644 --- a/cpp/include/cudf/dictionary/encode.hpp +++ b/cpp/include/cudf/dictionary/encode.hpp @@ -53,7 +53,7 @@ namespace dictionary { std::unique_ptr encode( column_view const& column, data_type indices_type = data_type{type_id::INT32}, - 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()); /** @@ -73,7 +73,7 @@ std::unique_ptr encode( */ std::unique_ptr decode( dictionary_column_view const& dictionary_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()); /** @} */ // end of group diff --git a/cpp/include/cudf/dictionary/search.hpp b/cpp/include/cudf/dictionary/search.hpp index d03c873f10fd..89ced89ebca1 100644 --- a/cpp/include/cudf/dictionary/search.hpp +++ b/cpp/include/cudf/dictionary/search.hpp @@ -36,7 +36,7 @@ namespace dictionary { std::unique_ptr get_index( dictionary_column_view const& dictionary, scalar const& key, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @} */ // end of group diff --git a/cpp/include/cudf/dictionary/update_keys.hpp b/cpp/include/cudf/dictionary/update_keys.hpp index bd48ae902b52..9a97567e7374 100644 --- a/cpp/include/cudf/dictionary/update_keys.hpp +++ b/cpp/include/cudf/dictionary/update_keys.hpp @@ -52,7 +52,7 @@ namespace dictionary { std::unique_ptr add_keys( dictionary_column_view const& dictionary_column, column_view const& new_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()); /** @@ -84,7 +84,7 @@ std::unique_ptr add_keys( std::unique_ptr remove_keys( dictionary_column_view const& dictionary_column, column_view const& keys_to_remove, - 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()); /** @@ -106,7 +106,7 @@ std::unique_ptr remove_keys( */ std::unique_ptr remove_unused_keys( dictionary_column_view const& dictionary_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()); /** @@ -143,7 +143,7 @@ std::unique_ptr remove_unused_keys( std::unique_ptr set_keys( dictionary_column_view const& dictionary_column, 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()); /** @@ -159,7 +159,7 @@ std::unique_ptr set_keys( */ std::vector> match_dictionaries( std::span 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/filling.hpp b/cpp/include/cudf/filling.hpp index b440dd43887e..aacc8ce06b1a 100644 --- a/cpp/include/cudf/filling.hpp +++ b/cpp/include/cudf/filling.hpp @@ -53,7 +53,7 @@ void fill_in_place(mutable_column_view& destination, size_type begin, size_type end, scalar const& value, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Fills a range of elements in a column out-of-place with a scalar @@ -83,7 +83,7 @@ std::unique_ptr fill( size_type begin, size_type end, scalar const& value, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -117,7 +117,7 @@ std::unique_ptr fill( std::unique_ptr
repeat( table_view const& input_table, column_view const& 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()); /** @@ -142,7 +142,7 @@ std::unique_ptr
repeat( std::unique_ptr
repeat( table_view const& input_table, 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()); /** @@ -174,7 +174,7 @@ std::unique_ptr sequence( size_type size, scalar const& init, scalar const& step, - 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()); /** @@ -202,7 +202,7 @@ std::unique_ptr sequence( std::unique_ptr sequence( size_type size, scalar const& init, - 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()); /** @@ -233,7 +233,7 @@ std::unique_ptr calendrical_month_sequence( size_type size, scalar const& init, size_type 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()); /** @} */ // end of group diff --git a/cpp/include/cudf/partitioning.hpp b/cpp/include/cudf/partitioning.hpp index 67cffce3e131..4b3a9b81dfaf 100644 --- a/cpp/include/cudf/partitioning.hpp +++ b/cpp/include/cudf/partitioning.hpp @@ -10,7 +10,7 @@ #include #include -#include +#include #include #include @@ -72,7 +72,7 @@ std::pair, std::vector> partition( table_view const& t, column_view const& partition_map, size_type num_partitions, - 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()); /** @@ -106,7 +106,7 @@ std::pair, std::vector> hash_partition( int num_partitions, hash_id hash_function = hash_id::HASH_MURMUR3, 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()); /** @@ -141,7 +141,7 @@ std::pair, std::vector> hash_partition( int num_partitions, hash_id hash_function = hash_id::HASH_MURMUR3, 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()); /** @@ -287,7 +287,7 @@ std::pair, std::vector> round_robi table_view const& input, cudf::size_type num_partitions, cudf::size_type start_partition = 0, - 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/scalar/scalar.hpp b/cpp/include/cudf/scalar/scalar.hpp index 12d35eec400d..41cf9d879302 100644 --- a/cpp/include/cudf/scalar/scalar.hpp +++ b/cpp/include/cudf/scalar/scalar.hpp @@ -12,10 +12,11 @@ #include #include -#include #include #include +#include + #include #include @@ -58,7 +59,7 @@ class scalar { * @param is_valid true: set the value to valid. false: set it to null. * @param stream CUDA stream used for device memory operations. */ - void set_valid_async(bool is_valid, rmm::cuda_stream_view stream = cudf::get_default_stream()); + void set_valid_async(bool is_valid, cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Indicates whether the scalar contains a valid value. @@ -70,7 +71,7 @@ class scalar { * @return true Value is valid * @return false Value is invalid/null */ - [[nodiscard]] bool is_valid(rmm::cuda_stream_view stream = cudf::get_default_stream()) const; + [[nodiscard]] bool is_valid(cuda::stream_ref stream = cudf::get_default_stream()) const; /** * @brief Returns a raw pointer to the validity bool in device memory. @@ -104,7 +105,7 @@ class scalar { * @param mr Device memory resource to use for device memory allocation. */ scalar(scalar 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()); /** @@ -120,7 +121,7 @@ class scalar { */ scalar(data_type type, bool is_valid = false, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); }; @@ -157,7 +158,7 @@ class fixed_width_scalar : public scalar { * @param mr Device memory resource to use for device memory allocation. */ fixed_width_scalar(fixed_width_scalar 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()); /** @@ -166,7 +167,7 @@ class fixed_width_scalar : public scalar { * @param value New value of scalar. * @param stream CUDA stream used for device memory operations. */ - void set_value(T value, rmm::cuda_stream_view stream = cudf::get_default_stream()); + void set_value(T value, cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Get the value of the scalar. @@ -174,7 +175,7 @@ class fixed_width_scalar : public scalar { * @param stream CUDA stream used for device memory operations. * @return Value of the scalar */ - [[nodiscard]] T value(rmm::cuda_stream_view stream = cudf::get_default_stream()) const; + [[nodiscard]] T value(cuda::stream_ref stream = cudf::get_default_stream()) const; /** * @brief Returns a raw pointer to the value in device memory. @@ -201,7 +202,7 @@ class fixed_width_scalar : public scalar { */ fixed_width_scalar(T value, bool is_valid = 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()); /** @@ -214,7 +215,7 @@ class fixed_width_scalar : public scalar { */ fixed_width_scalar(rmm::device_scalar&& data, bool is_valid = 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()); }; @@ -250,7 +251,7 @@ class numeric_scalar : public detail::fixed_width_scalar { * @param mr Device memory resource to use for device memory allocation. */ numeric_scalar(numeric_scalar 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()); /** @@ -263,7 +264,7 @@ class numeric_scalar : public detail::fixed_width_scalar { */ numeric_scalar(T value, bool is_valid = 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()); /** @@ -276,7 +277,7 @@ class numeric_scalar : public detail::fixed_width_scalar { */ numeric_scalar(rmm::device_scalar&& data, bool is_valid = 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()); }; @@ -313,7 +314,7 @@ class fixed_point_scalar : public scalar { * @param mr Device memory resource to use for device memory allocation. */ fixed_point_scalar(fixed_point_scalar 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()); /** @@ -328,7 +329,7 @@ class fixed_point_scalar : public scalar { fixed_point_scalar(rep_type value, numeric::scale_type scale, bool is_valid = 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()); /** @@ -341,7 +342,7 @@ class fixed_point_scalar : public scalar { */ fixed_point_scalar(rep_type value, bool is_valid = 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()); /** @@ -354,7 +355,7 @@ class fixed_point_scalar : public scalar { */ fixed_point_scalar(T value, bool is_valid = 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()); /** @@ -369,7 +370,7 @@ class fixed_point_scalar : public scalar { fixed_point_scalar(rmm::device_scalar&& data, numeric::scale_type scale, bool is_valid = 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()); /** @@ -378,7 +379,7 @@ class fixed_point_scalar : public scalar { * @param stream CUDA stream used for device memory operations. * @return The value of the scalar */ - [[nodiscard]] rep_type value(rmm::cuda_stream_view stream = cudf::get_default_stream()) const; + [[nodiscard]] rep_type value(cuda::stream_ref stream = cudf::get_default_stream()) const; /** * @brief Get the decimal32, decimal64 or decimal128. @@ -386,8 +387,7 @@ class fixed_point_scalar : public scalar { * @param stream CUDA stream used for device memory operations. * @return The decimal32, decimal64 or decimal128 value */ - [[nodiscard]] T fixed_point_value( - rmm::cuda_stream_view stream = cudf::get_default_stream()) const; + [[nodiscard]] T fixed_point_value(cuda::stream_ref stream = cudf::get_default_stream()) const; /** * @brief Returns a raw pointer to the value in device memory. @@ -433,7 +433,7 @@ class string_scalar : public scalar { * @param mr Device memory resource to use for device memory allocation. */ string_scalar(string_scalar 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()); /** @@ -448,7 +448,7 @@ class string_scalar : public scalar { */ string_scalar(std::string_view string, bool is_valid = 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()); /** @@ -463,7 +463,7 @@ class string_scalar : public scalar { */ string_scalar(value_type const& source, bool is_valid = 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()); /** @@ -478,7 +478,7 @@ class string_scalar : public scalar { */ string_scalar(rmm::device_scalar& data, bool is_valid = 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()); /** @@ -494,7 +494,7 @@ class string_scalar : public scalar { */ string_scalar(rmm::device_buffer&& data, bool is_valid = 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()); /** @@ -503,8 +503,7 @@ class string_scalar : public scalar { * @param stream CUDA stream used for device memory operations. * @return The value of the scalar in a host std::string */ - [[nodiscard]] std::string to_string( - rmm::cuda_stream_view stream = cudf::get_default_stream()) const; + [[nodiscard]] std::string to_string(cuda::stream_ref stream = cudf::get_default_stream()) const; /** * @brief Get the value of the scalar as a string_view. @@ -512,7 +511,7 @@ class string_scalar : public scalar { * @param stream CUDA stream used for device memory operations. * @return The value of the scalar as a string_view */ - [[nodiscard]] value_type value(rmm::cuda_stream_view stream = cudf::get_default_stream()) const; + [[nodiscard]] value_type value(cuda::stream_ref stream = cudf::get_default_stream()) const; /** * @brief Returns the size of the string in bytes. @@ -561,7 +560,7 @@ class chrono_scalar : public detail::fixed_width_scalar { * @param mr Device memory resource to use for device memory allocation. */ chrono_scalar(chrono_scalar 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()); /** @@ -574,7 +573,7 @@ class chrono_scalar : public detail::fixed_width_scalar { */ chrono_scalar(T value, bool is_valid = 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()); /** @@ -587,7 +586,7 @@ class chrono_scalar : public detail::fixed_width_scalar { */ chrono_scalar(rmm::device_scalar&& data, bool is_valid = 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()); }; @@ -620,7 +619,7 @@ class timestamp_scalar : public chrono_scalar { * @param mr Device memory resource to use for device memory allocation. */ timestamp_scalar(timestamp_scalar 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()); /** @@ -636,7 +635,7 @@ class timestamp_scalar : public chrono_scalar { template timestamp_scalar(Duration2 const& value, bool is_valid, - 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()); /** @@ -644,7 +643,7 @@ class timestamp_scalar : public chrono_scalar { * @param stream CUDA stream used for device memory operations. * @return The duration in number of ticks since the UNIX epoch */ - rep_type ticks_since_epoch(rmm::cuda_stream_view stream); + rep_type ticks_since_epoch(cuda::stream_ref stream); }; /** @@ -676,7 +675,7 @@ class duration_scalar : public chrono_scalar { * @param mr Device memory resource to use for device memory allocation. */ duration_scalar(duration_scalar 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()); /** @@ -689,7 +688,7 @@ class duration_scalar : public chrono_scalar { */ duration_scalar(rep_type value, bool is_valid, - 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()); /** @@ -697,7 +696,7 @@ class duration_scalar : public chrono_scalar { * @param stream CUDA stream used for device memory operations. * @return The duration in number of ticks */ - rep_type count(rmm::cuda_stream_view stream); + rep_type count(cuda::stream_ref stream); }; /** @@ -725,7 +724,7 @@ class list_scalar : public scalar { * @param mr Device memory resource to use for device memory allocation. */ list_scalar(list_scalar 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()); /** @@ -740,7 +739,7 @@ class list_scalar : public scalar { */ list_scalar(cudf::column_view const& data, bool is_valid = 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()); /** @@ -753,7 +752,7 @@ class list_scalar : public scalar { */ list_scalar(cudf::column&& data, bool is_valid = 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()); /** @@ -790,7 +789,7 @@ class struct_scalar : public scalar { * @param mr Device memory resource to use for device memory allocation. */ struct_scalar(struct_scalar 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()); /** @@ -805,7 +804,7 @@ class struct_scalar : public scalar { */ struct_scalar(table_view const& data, bool is_valid = 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()); /** @@ -820,7 +819,7 @@ class struct_scalar : public scalar { */ struct_scalar(std::span data, bool is_valid = 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()); /** @@ -836,7 +835,7 @@ class struct_scalar : public scalar { */ struct_scalar(table&& data, bool is_valid = 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()); /** @@ -864,7 +863,7 @@ class struct_scalar : public scalar { */ static table init_data(table&& data, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr); }; diff --git a/cpp/include/cudf/scalar/scalar_factories.hpp b/cpp/include/cudf/scalar/scalar_factories.hpp index 3433edd1a603..5a146c6dd879 100644 --- a/cpp/include/cudf/scalar/scalar_factories.hpp +++ b/cpp/include/cudf/scalar/scalar_factories.hpp @@ -10,6 +10,8 @@ #include +#include + #include /** @@ -37,7 +39,7 @@ namespace CUDF_EXPORT cudf { */ std::unique_ptr make_numeric_scalar( data_type type, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -54,7 +56,7 @@ std::unique_ptr make_numeric_scalar( */ std::unique_ptr make_timestamp_scalar( data_type type, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -71,7 +73,7 @@ std::unique_ptr make_timestamp_scalar( */ std::unique_ptr make_duration_scalar( data_type type, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -88,7 +90,7 @@ std::unique_ptr make_duration_scalar( */ std::unique_ptr make_fixed_width_scalar( data_type type, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -120,7 +122,7 @@ std::unique_ptr make_string_scalar( */ std::unique_ptr make_default_constructed_scalar( data_type type, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -135,7 +137,7 @@ std::unique_ptr make_default_constructed_scalar( */ std::unique_ptr make_empty_scalar_like( 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()); /** @@ -150,7 +152,7 @@ std::unique_ptr make_empty_scalar_like( template std::unique_ptr make_fixed_width_scalar( T value, - 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()) { return std::make_unique>(value, true, stream, mr); @@ -170,7 +172,7 @@ template std::unique_ptr make_fixed_point_scalar( typename T::rep value, numeric::scale_type scale, - 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()) { return std::make_unique>(value, scale, true, stream, mr); @@ -186,7 +188,7 @@ std::unique_ptr make_fixed_point_scalar( */ std::unique_ptr make_list_scalar( column_view elements, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); /** @@ -201,7 +203,7 @@ std::unique_ptr make_list_scalar( */ std::unique_ptr make_struct_scalar( table_view const& data, - 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()); /** @@ -216,7 +218,7 @@ std::unique_ptr make_struct_scalar( */ std::unique_ptr make_struct_scalar( std::span data, - 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/search.hpp b/cpp/include/cudf/search.hpp index 936b100f98fc..561499857c3c 100644 --- a/cpp/include/cudf/search.hpp +++ b/cpp/include/cudf/search.hpp @@ -64,7 +64,7 @@ std::unique_ptr lower_bound( table_view const& needles, 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()); /** @@ -106,7 +106,7 @@ std::unique_ptr upper_bound( table_view const& needles, 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()); /** @@ -129,7 +129,7 @@ std::unique_ptr upper_bound( */ bool contains(column_view const& haystack, scalar const& needle, - rmm::cuda_stream_view stream = cudf::get_default_stream()); + cuda::stream_ref stream = cudf::get_default_stream()); /** * @brief Check if the given `needles` values exists in the `haystack` column. @@ -155,7 +155,7 @@ bool contains(column_view const& haystack, std::unique_ptr contains( column_view const& haystack, column_view const& needles, - 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_test/column_utilities.hpp b/cpp/include/cudf_test/column_utilities.hpp index 57b3570be749..e26f64f19586 100644 --- a/cpp/include/cudf_test/column_utilities.hpp +++ b/cpp/include/cudf_test/column_utilities.hpp @@ -18,8 +18,7 @@ #include #include -#include - +#include #include #include @@ -57,7 +56,7 @@ bool expect_column_properties_equal( cudf::column_view const& lhs, cudf::column_view const& rhs, debug_output_level verbosity = debug_output_level::FIRST_ERROR, - rmm::cuda_stream_view stream = cudf::test::get_default_stream(), + cuda::stream_ref stream = cudf::test::get_default_stream(), cudf::memory_resources mr = cudf::get_current_device_resource_ref()); /** @@ -82,7 +81,7 @@ bool expect_column_properties_equivalent( cudf::column_view const& lhs, cudf::column_view const& rhs, debug_output_level verbosity = debug_output_level::FIRST_ERROR, - rmm::cuda_stream_view stream = cudf::test::get_default_stream(), + cuda::stream_ref stream = cudf::test::get_default_stream(), cudf::memory_resources mr = cudf::get_current_device_resource_ref()); /** @@ -104,7 +103,7 @@ bool expect_column_properties_equivalent( bool expect_columns_equal(cudf::column_view const& lhs, cudf::column_view const& rhs, debug_output_level verbosity = debug_output_level::FIRST_ERROR, - rmm::cuda_stream_view stream = cudf::test::get_default_stream(), + cuda::stream_ref stream = cudf::test::get_default_stream(), cudf::memory_resources mr = cudf::get_current_device_resource_ref()); /** @@ -130,7 +129,7 @@ bool expect_columns_equivalent(cudf::column_view const& lhs, cudf::column_view const& rhs, debug_output_level verbosity = debug_output_level::FIRST_ERROR, size_type fp_ulps = cudf::test::default_ulp, - rmm::cuda_stream_view stream = cudf::test::get_default_stream(), + cuda::stream_ref stream = cudf::test::get_default_stream(), cudf::memory_resources mr = cudf::get_current_device_resource_ref()); /** @@ -147,8 +146,8 @@ bool expect_columns_equivalent(cudf::column_view const& lhs, void expect_equal_buffers(void const* lhs, void const* rhs, std::size_t size_bytes, - rmm::cuda_stream_view stream = cudf::test::get_default_stream(), - cudf::memory_resources mr = cudf::get_current_device_resource_ref()); + cuda::stream_ref stream = cudf::test::get_default_stream(), + cudf::memory_resources mr = cudf::get_current_device_resource_ref()); } // namespace detail @@ -169,8 +168,8 @@ void expect_column_empty(cudf::column_view const& col); */ std::vector bitmask_to_host( cudf::column_view const& c, - rmm::cuda_stream_view stream = cudf::test::get_default_stream(), - cudf::memory_resources mr = cudf::get_current_device_resource_ref()); + cuda::stream_ref stream = cudf::test::get_default_stream(), + cudf::memory_resources mr = cudf::get_current_device_resource_ref()); /** * @brief Validates bitmask situated in host as per `number_of_elements` @@ -200,8 +199,8 @@ bool validate_host_masks(std::vector const& expected_mask, template ()>* = nullptr> std::pair, std::vector> to_host( column_view c, - rmm::cuda_stream_view stream = cudf::test::get_default_stream(), - cudf::memory_resources mr = cudf::get_current_device_resource_ref()) + cuda::stream_ref stream = cudf::test::get_default_stream(), + cudf::memory_resources mr = cudf::get_current_device_resource_ref()) { auto col_span = cudf::device_span(c.data(), c.size()); auto host_data = cudf::detail::make_host_vector(col_span, stream); @@ -227,8 +226,8 @@ std::pair, std::vector> to_host( template ()>* = nullptr> CUDF_EXPORT std::pair, std::vector> to_host( column_view c, - rmm::cuda_stream_view stream = cudf::test::get_default_stream(), - cudf::memory_resources mr = cudf::get_current_device_resource_ref()); + cuda::stream_ref stream = cudf::test::get_default_stream(), + cudf::memory_resources mr = cudf::get_current_device_resource_ref()); /** * @brief Copies the data and bitmask of a `column_view` of strings @@ -244,7 +243,7 @@ CUDF_EXPORT std::pair, std::vector> to_host */ template <> CUDF_EXPORT std::pair, std::vector> to_host( - column_view c, rmm::cuda_stream_view stream, cudf::memory_resources mr); + column_view c, cuda::stream_ref stream, cudf::memory_resources mr); //! @endcond /** diff --git a/cpp/src/column/column.cu b/cpp/src/column/column.cu index 5aeddc4f1063..085fe0d9562c 100644 --- a/cpp/src/column/column.cu +++ b/cpp/src/column/column.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 */ @@ -24,9 +24,9 @@ #include #include -#include #include +#include #include #include @@ -37,7 +37,7 @@ namespace cudf { // Copy ctor w/ optional stream/mr -column::column(column const& other, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) +column::column(column const& other, cuda::stream_ref stream, rmm::device_async_resource_ref mr) : _type{other._type}, _size{other._size}, _data{other._data, stream, mr}, @@ -137,7 +137,7 @@ void column::set_null_mask(rmm::device_buffer&& new_null_mask, size_type new_nul void column::set_null_mask(rmm::device_buffer const& new_null_mask, size_type new_null_count, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (new_null_count > 0) { CUDF_EXPECTS(new_null_mask.size() >= cudf::bitmask_allocation_size_bytes(this->size()), @@ -157,7 +157,7 @@ void column::set_null_count(size_type new_null_count) namespace { struct create_column_from_view { cudf::column_view view; - rmm::cuda_stream_view stream; + cuda::stream_ref stream; rmm::device_async_resource_ref mr; template @@ -253,14 +253,14 @@ struct create_column_from_view { } // anonymous namespace // Copy from a view -column::column(column_view view, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) +column::column(column_view view, cuda::stream_ref stream, rmm::device_async_resource_ref mr) : // Move is needed here because the dereference operator of unique_ptr returns // an lvalue reference, which would otherwise dispatch to the copy constructor column{std::move(*type_dispatcher(view.type(), create_column_from_view{view, stream, mr}))} { } -std::unique_ptr rebind_stream(column&& col, rmm::cuda_stream_view stream) +std::unique_ptr rebind_stream(column&& col, cuda::stream_ref stream) { auto const dtype = col.type(); auto const sz = col.size(); diff --git a/cpp/src/column/column_device_view.cu b/cpp/src/column/column_device_view.cu index 7698a14c886c..a46c0e442459 100644 --- a/cpp/src/column/column_device_view.cu +++ b/cpp/src/column/column_device_view.cu @@ -9,7 +9,7 @@ #include #include -#include +#include #include #include @@ -54,7 +54,7 @@ namespace { template std::unique_ptr> create_device_view_from_view(ColumnView const& source, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { size_type num_children = source.num_children(); @@ -112,7 +112,7 @@ column_device_view::column_device_view(column_view source, void* h_ptr, void* d_ // Construct a unique_ptr that invokes `destroy()` as it's deleter std::unique_ptr> column_device_view::create(column_view source, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { size_type num_children = source.num_children(); @@ -166,7 +166,7 @@ void mutable_column_device_view::destroy() { delete this; } // Construct a unique_ptr that invokes `destroy()` as it's deleter std::unique_ptr> mutable_column_device_view::create(mutable_column_view source, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return source.num_children() == 0 diff --git a/cpp/src/column/column_factories.cpp b/cpp/src/column/column_factories.cpp index 42867c4c6f1c..c7bb0d3d5471 100644 --- a/cpp/src/column/column_factories.cpp +++ b/cpp/src/column/column_factories.cpp @@ -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 */ @@ -62,7 +62,7 @@ std::unique_ptr make_empty_column(type_id id) { return make_empty_column std::unique_ptr make_numeric_column(data_type type, size_type size, mask_state state, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -84,7 +84,7 @@ std::unique_ptr make_numeric_column(data_type type, std::unique_ptr make_fixed_point_column(data_type type, size_type size, mask_state state, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -104,7 +104,7 @@ std::unique_ptr make_fixed_point_column(data_type type, std::unique_ptr make_timestamp_column(data_type type, size_type size, mask_state state, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -124,7 +124,7 @@ std::unique_ptr make_timestamp_column(data_type type, std::unique_ptr make_duration_column(data_type type, size_type size, mask_state state, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -144,7 +144,7 @@ std::unique_ptr make_duration_column(data_type type, std::unique_ptr make_fixed_width_column(data_type type, size_type size, mask_state state, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -162,7 +162,7 @@ std::unique_ptr make_fixed_width_column(data_type type, std::unique_ptr make_dictionary_from_scalar(scalar const& s, size_type size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (size == 0) return make_empty_column(type_id::DICTIONARY32); diff --git a/cpp/src/column/column_factories.cu b/cpp/src/column/column_factories.cu index 70d3b425aa69..bad8832b6a87 100644 --- a/cpp/src/column/column_factories.cu +++ b/cpp/src/column/column_factories.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 */ @@ -23,7 +23,7 @@ struct column_from_scalar_dispatch { template std::unique_ptr operator()(scalar const& value, size_type size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { if (size == 0) return make_empty_column(value.type()); @@ -41,7 +41,7 @@ template <> std::unique_ptr column_from_scalar_dispatch::operator()( scalar const& value, size_type size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { if (size == 0) return make_empty_column(value.type()); @@ -77,7 +77,7 @@ std::unique_ptr column_from_scalar_dispatch::operator() std::unique_ptr column_from_scalar_dispatch::operator()( - scalar const&, size_type, rmm::cuda_stream_view, rmm::device_async_resource_ref) const + scalar const&, size_type, cuda::stream_ref, rmm::device_async_resource_ref) const { CUDF_FAIL("dictionary not supported when creating from scalar"); } @@ -86,7 +86,7 @@ template <> std::unique_ptr column_from_scalar_dispatch::operator()( scalar const& value, size_type size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { auto lv = static_cast(&value); @@ -97,7 +97,7 @@ template <> std::unique_ptr column_from_scalar_dispatch::operator()( scalar const& value, size_type size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { CUDF_EXPECTS(size != 0, "0-length struct column is unsupported."); @@ -121,7 +121,7 @@ std::unique_ptr column_from_scalar_dispatch::operator() make_column_from_scalar(scalar const& s, size_type size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return type_dispatcher(s.type(), column_from_scalar_dispatch{}, s, size, stream, mr); diff --git a/cpp/src/column/column_view.cpp b/cpp/src/column/column_view.cpp index bcb619084e30..b9973b4a2eab 100644 --- a/cpp/src/column/column_view.cpp +++ b/cpp/src/column/column_view.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -135,7 +135,7 @@ column_view_base::column_view_base(data_type type, size_type column_view_base::null_count(size_type begin, size_type end, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { CUDF_EXPECTS((begin >= 0) && (end <= size()) && (begin <= end), "Range is out of bounds."); return (null_count() == 0) diff --git a/cpp/src/copying/concatenate.cu b/cpp/src/copying/concatenate.cu index f946c0df7fe5..4ebd84072da3 100644 --- a/cpp/src/copying/concatenate.cu +++ b/cpp/src/copying/concatenate.cu @@ -27,11 +27,11 @@ #include #include -#include #include #include #include +#include #include #include #include @@ -56,11 +56,11 @@ constexpr bool use_fused_kernel_heuristic(bool const has_nulls, size_t const num return has_nulls || num_columns > 4; } -auto create_device_views(host_span views, rmm::cuda_stream_view stream) +auto create_device_views(host_span views, cuda::stream_ref stream) { // Create device views for each input view using CDViewPtr = decltype(column_device_view::create(std::declval(), - std::declval())); + std::declval())); auto device_view_owners = std::vector(views.size()); std::transform(views.begin(), views.end(), device_view_owners.begin(), [stream](auto const& col) { return column_device_view::create(col, stream); @@ -153,14 +153,14 @@ size_type concatenate_masks(device_span d_views, device_span d_offsets, bitmask_type* dest_mask, size_type output_size, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { cudf::detail::device_scalar d_valid_count( 0, stream, cudf::get_current_device_resource_ref()); constexpr size_type block_size{256}; cudf::detail::grid_1d config(output_size, block_size); concatenate_masks_kernel - <<>>( + <<>>( d_views.data(), d_offsets.data(), static_cast(d_views.size()), @@ -173,7 +173,7 @@ size_type concatenate_masks(device_span d_views, size_type concatenate_masks(host_span views, bitmask_type* dest_mask, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // Preprocess and upload inputs to device memory auto const device_views = create_device_views(views, stream); @@ -239,7 +239,7 @@ CUDF_KERNEL void fused_concatenate_kernel(column_device_view const* input_views, template std::unique_ptr fused_concatenate(host_span views, bool const has_nulls, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { using mask_policy = cudf::mask_allocation_policy; @@ -268,7 +268,7 @@ std::unique_ptr fused_concatenate(host_span views, cudf::detail::grid_1d config(output_size, block_size); auto const kernel = has_nulls ? fused_concatenate_kernel : fused_concatenate_kernel; - kernel<<>>( + kernel<<>>( d_views.data(), d_offsets.data(), static_cast(d_views.size()), @@ -288,7 +288,7 @@ std::unique_ptr fused_concatenate(host_span views, template std::unique_ptr for_each_concatenate(host_span views, bool const has_nulls, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { size_type const total_element_count = @@ -329,7 +329,7 @@ std::unique_ptr for_each_concatenate(host_span views, struct concatenate_dispatch { host_span views; - rmm::cuda_stream_view stream; + cuda::stream_ref stream; rmm::device_async_resource_ref mr; // fixed width @@ -372,7 +372,7 @@ std::unique_ptr concatenate_dispatch::operator()() return cudf::structs::detail::concatenate(views, stream, mr); } -void bounds_and_type_check(host_span cols, rmm::cuda_stream_view stream); +void bounds_and_type_check(host_span cols, cuda::stream_ref stream); /** * @brief Functor for traversing child columns and recursively verifying concatenation @@ -382,7 +382,7 @@ class traverse_children { public: // nothing to do for simple types. template - void operator()(host_span, rmm::cuda_stream_view) + void operator()(host_span, cuda::stream_ref) { } @@ -407,7 +407,7 @@ class traverse_children { template <> void traverse_children::operator()(host_span cols, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // verify offsets check_offsets_size(cols); @@ -417,7 +417,7 @@ void traverse_children::operator()(host_span void traverse_children::operator()(host_span cols, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // march each child auto child_iter = cuda::counting_iterator{0}; @@ -440,7 +440,7 @@ void traverse_children::operator()(host_span void traverse_children::operator()(host_span cols, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // verify offsets check_offsets_size(cols); @@ -468,7 +468,7 @@ void traverse_children::operator()(host_span * * @throws cudf::logic_error if all of the input column types don't match */ -void bounds_and_type_check(host_span cols, rmm::cuda_stream_view stream) +void bounds_and_type_check(host_span cols, cuda::stream_ref stream) { // total size of all concatenated rows size_t const total_row_count = @@ -502,7 +502,7 @@ void bounds_and_type_check(host_span cols, rmm::cuda_stream_v // Concatenates the elements from a vector of column_views std::unique_ptr concatenate(std::span columns_to_concat, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(not columns_to_concat.empty(), "Unexpected empty list of columns to concatenate."); @@ -533,7 +533,7 @@ std::unique_ptr concatenate(std::span columns_to_conc } std::unique_ptr
concatenate(std::span tables_to_concat, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (tables_to_concat.empty()) { return std::make_unique
(); } @@ -576,7 +576,7 @@ std::unique_ptr
concatenate(std::span tables_to_concat, } rmm::device_buffer concatenate_masks(std::span views, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { bool const has_nulls = @@ -603,7 +603,7 @@ rmm::device_buffer concatenate_masks(std::span views, } // namespace detail rmm::device_buffer concatenate_masks(std::span views, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -612,7 +612,7 @@ rmm::device_buffer concatenate_masks(std::span views, // Concatenates the elements from a vector of column_views std::unique_ptr concatenate(std::span columns_to_concat, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -620,7 +620,7 @@ std::unique_ptr concatenate(std::span columns_to_conc } std::unique_ptr
concatenate(std::span tables_to_concat, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu index 5c0bfdd6d7a8..a30b6614dbf1 100644 --- a/cpp/src/copying/contiguous_split.cu +++ b/cpp/src/copying/contiguous_split.cu @@ -21,13 +21,13 @@ #include #include -#include #include #include #include #include #include +#include #include #include #include @@ -472,7 +472,7 @@ std::pair setup_source_buf_info(InputIter begin, InputIter end, src_buf_info* head, src_buf_info* current, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, int offset_stack_pos = 0, int parent_offset_index = -1, int offset_depth = 0); @@ -492,7 +492,7 @@ struct buf_info_functor { int offset_stack_pos, int parent_offset_index, int offset_depth, - rmm::cuda_stream_view) + cuda::stream_ref) { if (col.nullable()) { std::tie(current, offset_stack_pos) = @@ -535,7 +535,7 @@ std::pair buf_info_functor::operator() buf_info_functor::operator() buf_info_functor::operator() setup_source_buf_info(InputIter begin, InputIter end, src_buf_info* head, src_buf_info* current, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, int offset_stack_pos, int parent_offset_index, int offset_depth) @@ -1026,7 +1026,7 @@ struct packed_split_indices_and_src_buf_info { std::vector const& splits, std::size_t num_partitions, cudf::size_type num_src_bufs, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) : indices_size(cudf::util::round_up_safe((num_partitions + 1) * sizeof(int64_t), split_align)), src_buf_info_size( @@ -1087,7 +1087,7 @@ struct packed_split_indices_and_src_buf_info { struct packed_partition_buf_size_and_dst_buf_info { packed_partition_buf_size_and_dst_buf_info(std::size_t num_partitions, std::size_t num_bufs, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) : stream(stream), buf_sizes_size{cudf::util::round_up_safe(num_partitions * sizeof(std::size_t), split_align)}, @@ -1115,7 +1115,7 @@ struct packed_partition_buf_size_and_dst_buf_info { detail::cuda_memcpy_async(h_buf_sizes_and_dst_info, d_buf_sizes_and_dst_info, stream); } - rmm::cuda_stream_view const stream; + cuda::stream_ref const stream; // buffer sizes and destination info (used in batched copies) std::size_t const buf_sizes_size; @@ -1137,7 +1137,7 @@ struct packed_src_and_dst_pointers { packed_src_and_dst_pointers(cudf::table_view const& input, std::size_t num_partitions, cudf::size_type num_src_bufs, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) : stream(stream), src_bufs_size{cudf::util::round_up_safe(num_src_bufs * sizeof(uint8_t*), split_align)}, @@ -1166,7 +1166,7 @@ struct packed_src_and_dst_pointers { stream); } - rmm::cuda_stream_view const stream; + cuda::stream_ref const stream; std::size_t const src_bufs_size; std::size_t const dst_bufs_size; @@ -1199,7 +1199,7 @@ std::unique_ptr setup_src_and_dst_pointers( std::size_t num_partitions, cudf::size_type num_src_bufs, std::vector& out_buffers, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) { auto src_and_dst_pointers = std::make_unique( @@ -1236,7 +1236,7 @@ std::unique_ptr compute_splits( std::size_t num_partitions, cudf::size_type num_src_bufs, std::size_t num_bufs, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) { auto partition_buf_size_and_dst_buf_info = @@ -1367,7 +1367,7 @@ std::unique_ptr compute_splits( partition_buf_size_and_dst_buf_info->copy_to_host(); - stream.synchronize(); + stream.wait(); return partition_buf_size_and_dst_buf_info; } @@ -1384,7 +1384,7 @@ std::unique_ptr compute_splits( std::tuple> compute_num_bufs_and_splits(cudf::table_view const& input, std::vector const& splits, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) { std::size_t const num_partitions = splits.size() + 1; @@ -1435,7 +1435,7 @@ struct chunk_iteration_state { std::size_t const* const h_buf_sizes, std::size_t num_partitions, std::size_t user_buffer_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr); /** @@ -1495,7 +1495,7 @@ std::unique_ptr chunk_iteration_state::create( std::size_t const* const h_buf_sizes, std::size_t num_partitions, std::size_t user_buffer_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) { rmm::device_uvector d_batch_offsets(num_bufs + 1, stream, temp_mr); @@ -1702,7 +1702,7 @@ std::unique_ptr compute_batches(int num_bufs, std::size_t const* const h_buf_sizes, std::size_t num_partitions, std::size_t user_buffer_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) { // Since we parallelize at one block per copy, performance is vulnerable to situations where we @@ -1752,12 +1752,12 @@ void copy_data(int num_batches_to_copy, uint8_t** d_dst_bufs, device_span d_dst_buf_info, uint8_t* user_buffer, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { constexpr size_type block_size = 256; if (user_buffer != nullptr) { auto index_to_buffer = [user_buffer] __device__(unsigned int) { return user_buffer; }; - copy_partitions<<>>( + copy_partitions<<>>( index_to_buffer, d_src_bufs, d_dst_buf_info.data() + starting_batch); CUDF_CUDA_TRY(cudaGetLastError()); } else { @@ -1767,7 +1767,7 @@ void copy_data(int num_batches_to_copy, auto const dst_buf_index = dst_buf_info[buf_index].dst_buf_index; return d_dst_bufs[dst_buf_index]; }; - copy_partitions<<>>( + copy_partitions<<>>( index_to_buffer, d_src_bufs, d_dst_buf_info.data() + starting_batch); CUDF_CUDA_TRY(cudaGetLastError()); } @@ -1826,7 +1826,7 @@ namespace detail { struct contiguous_split_state { contiguous_split_state(cudf::table_view const& input, std::size_t user_buffer_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, std::optional mr, rmm::device_async_resource_ref temp_mr) : contiguous_split_state(input, {}, user_buffer_size, stream, mr, temp_mr) @@ -1835,7 +1835,7 @@ struct contiguous_split_state { contiguous_split_state(cudf::table_view const& input, std::vector const& splits, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, std::optional mr, rmm::device_async_resource_ref temp_mr) : contiguous_split_state(input, splits, 0, stream, mr, temp_mr) @@ -1953,7 +1953,7 @@ struct contiguous_split_state { contiguous_split_state(cudf::table_view const& input, std::vector const& splits, std::size_t user_buffer_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, std::optional mr, rmm::device_async_resource_ref temp_mr) : input(input), @@ -2094,7 +2094,7 @@ struct contiguous_split_state { cudf::table_view const input; ///< The input table_view to operate on std::size_t const user_buffer_size; ///< The size of the user buffer for the chunked_pack case - rmm::cuda_stream_view const stream; + cuda::stream_ref const stream; std::optional mr; ///< The resource for any data returned // this resource defaults to `mr` for the contiguous_split case, but it can be useful for the @@ -2145,7 +2145,7 @@ struct contiguous_split_state { std::vector contiguous_split(cudf::table_view const& input, std::vector const& splits, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // `temp_mr` is the same as `mr` for contiguous_split as it allocates all @@ -2159,7 +2159,7 @@ std::vector contiguous_split(cudf::table_view const& input, std::vector contiguous_split(cudf::table_view const& input, std::vector const& splits, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -2168,7 +2168,7 @@ std::vector contiguous_split(cudf::table_view const& input, chunked_pack::chunked_pack(cudf::table_view const& input, std::size_t user_buffer_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) { CUDF_EXPECTS(user_buffer_size >= desired_batch_size, @@ -2201,14 +2201,14 @@ std::unique_ptr> chunked_pack::build_metadata() const std::unique_ptr chunked_pack::create(cudf::table_view const& input, std::size_t user_buffer_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) { return std::make_unique(input, user_buffer_size, stream, temp_mr); } std::size_t packed_size(cudf::table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref temp_mr) { // Handle empty table cases diff --git a/cpp/src/copying/copy.cpp b/cpp/src/copying/copy.cpp index 3a286954c764..1ec3bb93ca0b 100644 --- a/cpp/src/copying/copy.cpp +++ b/cpp/src/copying/copy.cpp @@ -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 */ @@ -12,7 +12,7 @@ #include #include -#include +#include #include @@ -105,7 +105,7 @@ struct scalar_empty_like_functor { std::unique_ptr allocate_like(column_view const& input, size_type size, mask_allocation_policy mask_alloc, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -169,7 +169,7 @@ std::unique_ptr
empty_like(table_view const& input_table) std::unique_ptr allocate_like(column_view const& input, mask_allocation_policy mask_alloc, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -179,7 +179,7 @@ std::unique_ptr allocate_like(column_view const& input, std::unique_ptr allocate_like(column_view const& input, size_type size, mask_allocation_policy mask_alloc, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/copying/copy.cu b/cpp/src/copying/copy.cu index b6199ffe9b85..38d8c12bb4ad 100644 --- a/cpp/src/copying/copy.cu +++ b/cpp/src/copying/copy.cu @@ -19,11 +19,11 @@ #include #include -#include #include #include #include +#include #include @@ -45,14 +45,14 @@ struct copy_if_else_functor_impl { */ struct get_iterable_device_view { template - auto operator()(T const& input, rmm::cuda_stream_view stream) + auto operator()(T const& input, cuda::stream_ref stream) requires(std::is_same_v) { return cudf::column_device_view::create(input, stream); } template - auto operator()(T const& input, rmm::cuda_stream_view) + auto operator()(T const& input, cuda::stream_ref) requires(std::is_same_v) { return &input; @@ -68,7 +68,7 @@ struct copy_if_else_functor_impl bool left_nullable, bool right_nullable, Filter filter, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto p_lhs = get_iterable_device_view{}(lhs_h, stream); @@ -101,7 +101,7 @@ struct copy_if_else_functor_impl { bool left_nullable, bool right_nullable, Filter filter, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { using T = string_view; @@ -153,7 +153,7 @@ std::unique_ptr scatter_gather_based_if_else(cudf::column_view const& lh cudf::column_view const& rhs, size_type size, Filter is_left, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto gather_map = rmm::device_uvector{static_cast(size), stream}; @@ -187,7 +187,7 @@ std::unique_ptr scatter_gather_based_if_else(cudf::scalar const& lhs, cudf::column_view const& rhs, size_type size, Filter is_left, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto scatter_map = rmm::device_uvector{static_cast(size), stream}; @@ -216,7 +216,7 @@ std::unique_ptr scatter_gather_based_if_else(cudf::column_view const& lh cudf::scalar const& rhs, size_type size, Filter is_left, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return scatter_gather_based_if_else(rhs, lhs, size, logical_not{is_left}, stream, mr); @@ -227,7 +227,7 @@ std::unique_ptr scatter_gather_based_if_else(cudf::scalar const& lhs, cudf::scalar const& rhs, size_type size, Filter is_left, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto rhs_col = cudf::make_column_from_scalar(rhs, size, stream, mr); @@ -243,7 +243,7 @@ struct copy_if_else_functor_impl { bool, bool, Filter filter, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return scatter_gather_based_if_else(lhs, rhs, size, filter, stream, mr); @@ -259,7 +259,7 @@ struct copy_if_else_functor_impl { bool, bool, Filter filter, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return scatter_gather_based_if_else(lhs, rhs, size, filter, stream, mr); @@ -275,7 +275,7 @@ struct copy_if_else_functor_impl { bool, bool, Filter filter, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return scatter_gather_based_if_else(lhs, rhs, size, filter, stream, mr); @@ -294,7 +294,7 @@ struct copy_if_else_functor { bool left_nullable, bool right_nullable, Filter filter, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { copy_if_else_functor_impl copier{}; @@ -309,7 +309,7 @@ std::unique_ptr copy_if_else(Left const& lhs, bool left_nullable, bool right_nullable, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(boolean_mask.type() == data_type(type_id::BOOL8), @@ -347,7 +347,7 @@ std::unique_ptr copy_if_else(Left const& lhs, std::unique_ptr copy_if_else(column_view const& lhs, column_view const& rhs, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(boolean_mask.size() == lhs.size(), @@ -364,7 +364,7 @@ std::unique_ptr copy_if_else(column_view const& lhs, std::unique_ptr copy_if_else(scalar const& lhs, column_view const& rhs, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(boolean_mask.size() == rhs.size(), @@ -379,7 +379,7 @@ std::unique_ptr copy_if_else(scalar const& lhs, std::unique_ptr copy_if_else(column_view const& lhs, scalar const& rhs, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(boolean_mask.size() == lhs.size(), @@ -394,7 +394,7 @@ std::unique_ptr copy_if_else(column_view const& lhs, std::unique_ptr copy_if_else(scalar const& lhs, scalar const& rhs, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS( @@ -408,7 +408,7 @@ std::unique_ptr copy_if_else(scalar const& lhs, std::unique_ptr copy_if_else(column_view const& lhs, column_view const& rhs, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -418,7 +418,7 @@ std::unique_ptr copy_if_else(column_view const& lhs, std::unique_ptr copy_if_else(scalar const& lhs, column_view const& rhs, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -428,7 +428,7 @@ std::unique_ptr copy_if_else(scalar const& lhs, std::unique_ptr copy_if_else(column_view const& lhs, scalar const& rhs, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -438,7 +438,7 @@ std::unique_ptr copy_if_else(column_view const& lhs, std::unique_ptr copy_if_else(scalar const& lhs, scalar const& rhs, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/copying/copy_range.cu b/cpp/src/copying/copy_range.cu index 76f449193b0f..a84ae741960b 100644 --- a/cpp/src/copying/copy_range.cu +++ b/cpp/src/copying/copy_range.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 */ @@ -24,9 +24,8 @@ #include #include -#include - #include +#include #include #include @@ -38,7 +37,7 @@ void in_place_copy_range(cudf::column_view const& source, cudf::size_type source_begin, cudf::size_type source_end, cudf::size_type target_begin, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto p_source_device_view = cudf::column_device_view::create(source, stream); if (source.has_nulls()) { @@ -67,7 +66,7 @@ struct in_place_copy_range_dispatch { void operator()(cudf::size_type source_begin, cudf::size_type source_end, cudf::size_type target_begin, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { in_place_copy_range(source, target, source_begin, source_end, target_begin, stream); } @@ -88,7 +87,7 @@ struct out_of_place_copy_range_dispatch { cudf::size_type source_begin, cudf::size_type source_end, cudf::size_type target_begin, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()) { auto p_ret = std::make_unique(target, stream, mr); @@ -119,7 +118,7 @@ std::unique_ptr out_of_place_copy_range_dispatch::operator() out_of_place_copy_range_dispatch::operator() copy_range(column_view const& source, size_type source_begin, size_type source_end, size_type target_begin, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS((source_begin >= 0) && (source_end <= source.size()) && @@ -249,7 +248,7 @@ void copy_range_in_place(column_view const& source, size_type source_begin, size_type source_end, size_type target_begin, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); return detail::copy_range_in_place( @@ -261,7 +260,7 @@ std::unique_ptr copy_range(column_view const& source, size_type source_begin, size_type source_end, size_type target_begin, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/copying/gather.cu b/cpp/src/copying/gather.cu index 5a3288646b42..a43e662478a0 100644 --- a/cpp/src/copying/gather.cu +++ b/cpp/src/copying/gather.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 */ @@ -14,9 +14,8 @@ #include #include -#include - #include +#include #include #include @@ -28,7 +27,7 @@ std::unique_ptr
gather(table_view const& source_table, column_view const& gather_map, out_of_bounds_policy bounds_policy, negative_index_policy neg_indices, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(not gather_map.has_nulls(), "gather_map contains nulls", std::invalid_argument); @@ -55,7 +54,7 @@ std::unique_ptr
gather(table_view const& source_table, device_span const gather_map, out_of_bounds_policy bounds_policy, negative_index_policy neg_indices, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(gather_map.size() <= static_cast(std::numeric_limits::max()), @@ -74,7 +73,7 @@ std::unique_ptr
gather(table_view const& source_table, std::unique_ptr
gather(table_view const& source_table, column_view const& gather_map, out_of_bounds_policy bounds_policy, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -89,7 +88,7 @@ std::unique_ptr
gather(table_view const& source_table, column_view const& gather_map, out_of_bounds_policy bounds_policy, negative_index_policy neg_indices, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/copying/get_element.cu b/cpp/src/copying/get_element.cu index 1d989f7c9fb5..3ff512d15465 100644 --- a/cpp/src/copying/get_element.cu +++ b/cpp/src/copying/get_element.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 */ @@ -19,7 +19,7 @@ #include #include -#include +#include #include @@ -32,7 +32,7 @@ struct get_element_functor { template () && !is_fixed_point()>* p = nullptr> std::unique_ptr operator()(column_view const& input, size_type index, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto s = make_fixed_width_scalar(data_type(type_to_id()), stream, mr); @@ -55,7 +55,7 @@ struct get_element_functor { template >* p = nullptr> std::unique_ptr operator()(column_view const& input, size_type index, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto device_col = column_device_view::create(input, stream); @@ -79,7 +79,7 @@ struct get_element_functor { template >* p = nullptr> std::unique_ptr operator()(column_view const& input, size_type index, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto dict_view = dictionary_column_view(input); @@ -113,7 +113,7 @@ struct get_element_functor { template >* p = nullptr> std::unique_ptr operator()(column_view const& input, size_type index, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { bool valid = is_element_valid_sync(input, index, stream); @@ -137,7 +137,7 @@ struct get_element_functor { template ()>* p = nullptr> std::unique_ptr operator()(column_view const& input, size_type index, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { using Type = typename T::rep; @@ -163,7 +163,7 @@ struct get_element_functor { template >* p = nullptr> std::unique_ptr operator()(column_view const& input, size_type index, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { bool valid = is_element_valid_sync(input, index, stream); @@ -178,7 +178,7 @@ struct get_element_functor { std::unique_ptr get_element(column_view const& input, size_type index, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(index >= 0 and index < input.size(), "Index out of bounds", std::out_of_range); @@ -189,7 +189,7 @@ std::unique_ptr get_element(column_view const& input, std::unique_ptr get_element(column_view const& input, size_type index, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/copying/pack.cpp b/cpp/src/copying/pack.cpp index 58785aeda761..9cb3dc29ca30 100644 --- a/cpp/src/copying/pack.cpp +++ b/cpp/src/copying/pack.cpp @@ -7,7 +7,7 @@ #include #include -#include +#include #include #include @@ -272,7 +272,7 @@ table_view unpack(uint8_t const* metadata, uint8_t const* gpu_data) * @copydoc cudf::detail::pack */ packed_columns pack(cudf::table_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // do a contiguous_split with no splits to get the memory for the table @@ -441,7 +441,7 @@ packed_metadata_view::column_view packed_metadata_view::column(size_type i) cons * @copydoc cudf::pack */ packed_columns pack(cudf::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/copying/purge_nonempty_nulls.cu b/cpp/src/copying/purge_nonempty_nulls.cu index 8198c6942b29..e8a3299bf254 100644 --- a/cpp/src/copying/purge_nonempty_nulls.cu +++ b/cpp/src/copying/purge_nonempty_nulls.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 */ #include @@ -26,7 +26,7 @@ bool type_may_have_nonempty_nulls(cudf::type_id const& type) } /// Check if the (STRING/LIST) column has any null rows with non-zero length. -bool has_nonempty_null_rows(cudf::column_view const& input, rmm::cuda_stream_view stream) +bool has_nonempty_null_rows(cudf::column_view const& input, cuda::stream_ref stream) { if (not input.has_nulls()) { return false; } // No nulls => no dirty rows. @@ -53,7 +53,7 @@ bool has_nonempty_null_rows(cudf::column_view const& input, rmm::cuda_stream_vie /** * @copydoc cudf::detail::has_nonempty_nulls */ -bool has_nonempty_nulls(cudf::column_view const& input, rmm::cuda_stream_view stream) +bool has_nonempty_nulls(cudf::column_view const& input, cuda::stream_ref stream) { auto const type = input.type().id(); @@ -77,7 +77,7 @@ bool has_nonempty_nulls(cudf::column_view const& input, rmm::cuda_stream_view st } std::unique_ptr purge_nonempty_nulls(column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // If not compound types (LIST/STRING/STRUCT/DICTIONARY) then just copy the input into output. @@ -117,7 +117,7 @@ bool may_have_nonempty_nulls(column_view const& input) /** * @copydoc cudf::has_nonempty_nulls */ -bool has_nonempty_nulls(column_view const& input, rmm::cuda_stream_view stream) +bool has_nonempty_nulls(column_view const& input, cuda::stream_ref stream) { return detail::has_nonempty_nulls(input, stream); } @@ -126,7 +126,7 @@ bool has_nonempty_nulls(column_view const& input, rmm::cuda_stream_view stream) * @copydoc cudf::purge_nonempty_nulls(column_view const&, rmm::device_async_resource_ref) */ std::unique_ptr purge_nonempty_nulls(column_view const& input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return detail::purge_nonempty_nulls(input, stream, mr); diff --git a/cpp/src/copying/reverse.cu b/cpp/src/copying/reverse.cu index f394bec86c72..e0dae162d928 100644 --- a/cpp/src/copying/reverse.cu +++ b/cpp/src/copying/reverse.cu @@ -12,17 +12,17 @@ #include #include -#include #include #include #include +#include #include namespace cudf { namespace detail { std::unique_ptr
reverse(table_view const& source_table, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { size_type num_rows = source_table.num_rows(); @@ -36,7 +36,7 @@ std::unique_ptr
reverse(table_view const& source_table, } std::unique_ptr reverse(column_view const& source_column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return std::move( @@ -45,7 +45,7 @@ std::unique_ptr reverse(column_view const& source_column, } // namespace detail std::unique_ptr
reverse(table_view const& source_table, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -53,7 +53,7 @@ std::unique_ptr
reverse(table_view const& source_table, } std::unique_ptr reverse(column_view const& source_column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/copying/sample.cu b/cpp/src/copying/sample.cu index 26b3c174c07b..31e1ccba0000 100644 --- a/cpp/src/copying/sample.cu +++ b/cpp/src/copying/sample.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,9 @@ #include #include -#include - #include #include +#include #include #include #include @@ -30,7 +29,7 @@ std::unique_ptr
sample(table_view const& input, size_type const n, sample_with_replacement replacement, int64_t const seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(n >= 0, "expected number of samples should be non-negative"); @@ -85,7 +84,7 @@ std::unique_ptr
sample(table_view const& input, size_type const n, sample_with_replacement replacement, int64_t const seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/copying/scatter.cu b/cpp/src/copying/scatter.cu index df790e6905bb..c8b10f15740c 100644 --- a/cpp/src/copying/scatter.cu +++ b/cpp/src/copying/scatter.cu @@ -26,10 +26,9 @@ #include #include -#include - #include #include +#include #include #include #include @@ -67,7 +66,7 @@ void scatter_scalar_bitmask_inplace(std::reference_wrapper const& MapIterator scatter_map, size_type num_scatter_rows, column& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { constexpr size_type block_size = 256; @@ -85,7 +84,7 @@ void scatter_scalar_bitmask_inplace(std::reference_wrapper const& auto bitmask_kernel = source_is_valid ? marking_bitmask_kernel : marking_bitmask_kernel; - bitmask_kernel<<>>( + bitmask_kernel<<>>( *target_view, scatter_map, num_scatter_rows); CUDF_CUDA_TRY(cudaGetLastError()); @@ -100,7 +99,7 @@ struct column_scalar_scatterer_impl { MapIterator scatter_iter, size_type scatter_rows, column_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { CUDF_EXPECTS(cudf::have_same_types(target, source.get()), @@ -133,7 +132,7 @@ struct column_scalar_scatterer_impl { MapIterator scatter_iter, size_type scatter_rows, column_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { CUDF_EXPECTS(cudf::have_same_types(target, source.get()), @@ -157,7 +156,7 @@ struct column_scalar_scatterer_impl { MapIterator scatter_iter, size_type scatter_rows, column_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { CUDF_EXPECTS(source.get().type() == target.type(), @@ -177,7 +176,7 @@ struct column_scalar_scatterer_impl { MapIterator scatter_iter, size_type scatter_rows, column_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { auto dict_target = dictionary::detail::add_keys( @@ -218,7 +217,7 @@ struct column_scalar_scatterer { MapIterator scatter_iter, size_type scatter_rows, column_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { column_scalar_scatterer_impl scatterer{}; @@ -232,7 +231,7 @@ struct column_scalar_scatterer_impl { MapIterator scatter_iter, size_type scatter_rows, column_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { CUDF_EXPECTS(source.get().type() == target.type(), @@ -286,7 +285,7 @@ struct column_scalar_scatterer_impl { std::unique_ptr
scatter(table_view const& source, column_view const& scatter_map, table_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(source.num_columns() == target.num_columns(), @@ -311,7 +310,7 @@ std::unique_ptr
scatter(table_view const& source, std::unique_ptr
scatter(table_view const& source, device_span const scatter_map, table_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(scatter_map.size() <= static_cast(std::numeric_limits::max()), @@ -328,7 +327,7 @@ std::unique_ptr
scatter(table_view const& source, std::unique_ptr
scatter(std::vector> const& source, column_view const& indices, table_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(source.size() == static_cast(target.num_columns()), @@ -380,7 +379,7 @@ std::unique_ptr
scatter(std::vector> std::unique_ptr boolean_mask_scatter(column_view const& input, column_view const& target, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto indices = cudf::make_numeric_column(data_type{type_id::INT32}, @@ -411,7 +410,7 @@ std::unique_ptr boolean_mask_scatter(column_view const& input, std::unique_ptr boolean_mask_scatter(scalar const& input, column_view const& target, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return detail::copy_if_else(input, target, boolean_mask, stream, mr); @@ -420,7 +419,7 @@ std::unique_ptr boolean_mask_scatter(scalar const& input, std::unique_ptr
boolean_mask_scatter(table_view const& input, table_view const& target, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input.num_columns() == target.num_columns(), @@ -462,7 +461,7 @@ std::unique_ptr
boolean_mask_scatter( std::vector> const& input, table_view const& target, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(static_cast(input.size()) == target.num_columns(), @@ -506,7 +505,7 @@ std::unique_ptr
boolean_mask_scatter( std::unique_ptr
scatter(table_view const& source, column_view const& scatter_map, table_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -516,7 +515,7 @@ std::unique_ptr
scatter(table_view const& source, std::unique_ptr
scatter(std::vector> const& source, column_view const& indices, table_view const& target, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -526,7 +525,7 @@ std::unique_ptr
scatter(std::vector> std::unique_ptr
boolean_mask_scatter(table_view const& input, table_view const& target, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -537,7 +536,7 @@ std::unique_ptr
boolean_mask_scatter( std::vector> const& input, table_view const& target, column_view const& boolean_mask, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/copying/segmented_shift.cu b/cpp/src/copying/segmented_shift.cu index 2baf35ff0460..0b6d52f8f2d0 100644 --- a/cpp/src/copying/segmented_shift.cu +++ b/cpp/src/copying/segmented_shift.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -13,8 +13,7 @@ #include #include -#include - +#include #include #include #include @@ -62,7 +61,7 @@ struct segmented_shift_functor() device_span segment_offsets, size_type offset, scalar const& fill_value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto values_device_view = column_device_view::create(segmented_values, stream); @@ -91,7 +90,7 @@ struct segmented_shift_functor { device_span segment_offsets, size_type offset, scalar const& fill_value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto values_device_view = column_device_view::create(segmented_values, stream); @@ -118,7 +117,7 @@ struct segmented_shift_functor_forwarder { device_span segment_offsets, size_type offset, scalar const& fill_value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { segmented_shift_functor shifter; @@ -132,7 +131,7 @@ std::unique_ptr segmented_shift(column_view const& segmented_values, device_span segment_offsets, size_type offset, scalar const& fill_value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (segmented_values.is_empty()) { return empty_like(segmented_values); } diff --git a/cpp/src/copying/shift.cu b/cpp/src/copying/shift.cu index ea818e8633ce..2a8834078505 100644 --- a/cpp/src/copying/shift.cu +++ b/cpp/src/copying/shift.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 */ @@ -19,10 +19,10 @@ #include #include -#include #include #include +#include #include #include @@ -41,7 +41,7 @@ inline bool __device__ out_of_bounds(size_type size, size_type idx) std::pair create_null_mask(column_device_view const& input, size_type offset, scalar const& fill_value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const size = input.size(); @@ -69,7 +69,7 @@ struct shift_functor { std::unique_ptr operator()(column_view const& input, size_type offset, scalar const& fill_value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(std::is_same_v) { @@ -89,7 +89,7 @@ struct shift_functor { std::unique_ptr operator()(column_view const& input, size_type offset, scalar const& fill_value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(cudf::is_fixed_width()) { @@ -148,7 +148,7 @@ namespace detail { std::unique_ptr shift(column_view const& input, size_type offset, scalar const& fill_value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(cudf::have_same_types(input, fill_value), @@ -166,7 +166,7 @@ std::unique_ptr shift(column_view const& input, std::unique_ptr shift(column_view const& input, size_type offset, scalar const& fill_value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/copying/slice.cu b/cpp/src/copying/slice.cu index 9c63022cd307..a08d751ea33d 100644 --- a/cpp/src/copying/slice.cu +++ b/cpp/src/copying/slice.cu @@ -13,8 +13,7 @@ #include #include -#include - +#include #include #include @@ -24,10 +23,7 @@ namespace cudf { namespace detail { template -ColumnView slice(ColumnView const& input, - size_type begin, - size_type end, - rmm::cuda_stream_view stream) +ColumnView slice(ColumnView const& input, size_type begin, size_type end, cuda::stream_ref stream) { CUDF_EXPECTS(begin >= 0, "Invalid beginning of range.", std::out_of_range); CUDF_EXPECTS(end >= begin, "Invalid end of range.", std::invalid_argument); @@ -49,18 +45,15 @@ ColumnView slice(ColumnView const& input, children); } -template column_view slice(column_view const&, - size_type, - size_type, - rmm::cuda_stream_view); +template column_view slice(column_view const&, size_type, size_type, cuda::stream_ref); template mutable_column_view slice(mutable_column_view const&, size_type, size_type, - rmm::cuda_stream_view); + cuda::stream_ref); std::vector slice(column_view const& input, std::span indices, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS(indices.size() % 2 == 0, "indices size must be even", std::invalid_argument); @@ -96,7 +89,7 @@ std::vector slice(column_view const& input, std::vector slice(table_view const& input, std::span indices, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS(indices.size() % 2 == 0, "indices size must be even", std::invalid_argument); if (indices.empty()) { return {}; } @@ -141,14 +134,14 @@ std::vector slice(table_view const& input, std::vector slice(column_view const& input, std::initializer_list indices, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { return detail::slice(input, std::span(indices.begin(), indices.size()), stream); } std::vector slice(table_view const& input, std::initializer_list indices, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { return detail::slice(input, std::span(indices.begin(), indices.size()), stream); }; @@ -157,7 +150,7 @@ std::vector slice(table_view const& input, std::vector slice(column_view const& input, std::span indices, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); return detail::slice(input, indices, stream); @@ -165,7 +158,7 @@ std::vector slice(column_view const& input, std::vector slice(table_view const& input, std::span indices, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); return detail::slice(input, indices, stream); @@ -173,7 +166,7 @@ std::vector slice(table_view const& input, std::vector slice(column_view const& input, std::initializer_list indices, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); return detail::slice(input, indices, stream); @@ -181,7 +174,7 @@ std::vector slice(column_view const& input, std::vector slice(table_view const& input, std::initializer_list indices, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); return detail::slice(input, indices, stream); diff --git a/cpp/src/copying/split.cpp b/cpp/src/copying/split.cpp index 74f9bf0d9bb5..a75574f543d1 100644 --- a/cpp/src/copying/split.cpp +++ b/cpp/src/copying/split.cpp @@ -7,7 +7,7 @@ #include #include -#include +#include #include #include @@ -19,7 +19,7 @@ template std::vector split(T const& input, size_type column_size, std::span splits, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (splits.empty() or column_size == 0) { return std::vector{input}; } CUDF_EXPECTS( @@ -41,14 +41,14 @@ std::vector split(T const& input, std::vector split(cudf::column_view const& input, std::span splits, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { return split(input, input.size(), splits, stream); } std::vector split(cudf::table_view const& input, std::span splits, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // A genuinely empty table (no columns and no rows) has nothing to split. if (input.num_columns() == 0 && input.num_rows() == 0) { return {}; } @@ -57,14 +57,14 @@ std::vector split(cudf::table_view const& input, std::vector split(column_view const& input, std::initializer_list splits, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { return detail::split(input, std::span(splits.begin(), splits.size()), stream); } std::vector split(table_view const& input, std::initializer_list splits, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { return detail::split(input, std::span(splits.begin(), splits.size()), stream); } @@ -73,7 +73,7 @@ std::vector split(table_view const& input, std::vector split(cudf::column_view const& input, std::span splits, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); return detail::split(input, splits, stream); @@ -81,7 +81,7 @@ std::vector split(cudf::column_view const& input, std::vector split(cudf::table_view const& input, std::span splits, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); return detail::split(input, splits, stream); @@ -89,7 +89,7 @@ std::vector split(cudf::table_view const& input, std::vector split(column_view const& input, std::initializer_list splits, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); return detail::split(input, splits, stream); @@ -97,7 +97,7 @@ std::vector split(column_view const& input, std::vector split(table_view const& input, std::initializer_list splits, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); return detail::split(input, splits, stream); diff --git a/cpp/src/dictionary/add_keys.cu b/cpp/src/dictionary/add_keys.cu index 39ccc27e93e3..017ca6e86301 100644 --- a/cpp/src/dictionary/add_keys.cu +++ b/cpp/src/dictionary/add_keys.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 */ @@ -24,7 +24,7 @@ namespace detail { std::unique_ptr add_keys(dictionary_column_view const& input, column_view const& new_keys, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(!new_keys.has_nulls(), "Keys must not have nulls", std::invalid_argument); @@ -59,7 +59,7 @@ std::unique_ptr add_keys(dictionary_column_view const& input, std::unique_ptr add_keys(dictionary_column_view const& dictionary_column, column_view 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/dictionary/decode.cu b/cpp/src/dictionary/decode.cu index 6265815b7487..0338a939ed7d 100644 --- a/cpp/src/dictionary/decode.cu +++ b/cpp/src/dictionary/decode.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -14,7 +14,7 @@ #include #include -#include +#include namespace cudf { namespace dictionary { @@ -35,7 +35,7 @@ struct indices_handler_fn { * @brief Decode a column from a dictionary. */ std::unique_ptr decode(dictionary_column_view const& source, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (source.is_empty()) return make_empty_column(type_id::EMPTY); @@ -66,7 +66,7 @@ std::unique_ptr decode(dictionary_column_view const& source, } // namespace detail std::unique_ptr decode(dictionary_column_view const& source, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/dictionary/detail/concatenate.cu b/cpp/src/dictionary/detail/concatenate.cu index 9b2b275aa9c9..e3e428c14ace 100644 --- a/cpp/src/dictionary/detail/concatenate.cu +++ b/cpp/src/dictionary/detail/concatenate.cu @@ -21,7 +21,6 @@ #include #include -#include #include #include #include @@ -30,6 +29,7 @@ #include #include #include +#include #include #include #include @@ -105,7 +105,7 @@ struct compute_children_offsets_fn { * @param stream Stream used for allocating the output rmm::device_uvector. * @return Vector of offsets_pair objects for keys and indices. */ - rmm::device_uvector create_children_offsets(rmm::cuda_stream_view stream) + rmm::device_uvector create_children_offsets(cuda::stream_ref stream) { auto offsets = cudf::detail::make_host_vector(columns_ptrs.size(), stream); thrust::transform_exclusive_scan( @@ -154,7 +154,7 @@ struct map_indices_fn { } // namespace std::unique_ptr concatenate(host_span columns, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // exception here is the same behavior as in cudf::concatenate @@ -196,7 +196,7 @@ std::unique_ptr concatenate(host_span columns, auto probe = encode_probe_t{row_hash.device_hasher(cudf::nullate::NO{})}; auto allocator = rmm::mr::polymorphic_allocator(cudf::get_current_device_resource_ref()); auto set = cuco::static_set{ - all_keys->size(), 0.5, empty_key, d_equal, probe, {}, {}, allocator, stream.value()}; + all_keys->size(), 0.5, empty_key, d_equal, probe, {}, {}, allocator, stream.get()}; auto set_ref = set.ref(cuco::insert_and_find); using set_ref_t = decltype(set_ref); @@ -210,7 +210,7 @@ std::unique_ptr concatenate(host_span columns, policy, iota, iota + all_keys->size(), d_indices.begin(), insert_keys_fn{set_ref, *d_all_keys}); auto keys_indices = rmm::device_uvector( all_keys->size(), stream, cudf::get_current_device_resource_ref()); - auto keys_end = set.retrieve_all(keys_indices.begin(), stream.value()); + auto keys_end = set.retrieve_all(keys_indices.begin(), stream.get()); keys_indices.resize(cuda::std::distance(keys_indices.begin(), keys_end), stream); // use keys_indices to retrieve the keys (gather) diff --git a/cpp/src/dictionary/detail/merge.cu b/cpp/src/dictionary/detail/merge.cu index b9968e46348d..d946c68bcdf8 100644 --- a/cpp/src/dictionary/detail/merge.cu +++ b/cpp/src/dictionary/detail/merge.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 */ @@ -13,9 +13,9 @@ #include #include -#include #include +#include #include namespace cudf { @@ -25,7 +25,7 @@ namespace detail { std::unique_ptr merge(dictionary_column_view const& lcol, dictionary_column_view const& rcol, cudf::detail::index_vector const& row_order, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const lcol_iter = cudf::detail::indexalator_factory::make_input_iterator(lcol.indices()); diff --git a/cpp/src/dictionary/dictionary_factories.cu b/cpp/src/dictionary/dictionary_factories.cu index 064ea1ce49b5..05df8d36b3fd 100644 --- a/cpp/src/dictionary/dictionary_factories.cu +++ b/cpp/src/dictionary/dictionary_factories.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 */ @@ -12,14 +12,14 @@ #include #include -#include +#include namespace cudf { namespace { struct dispatch_create_indices { template std::unique_ptr operator()(column_view const& indices, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(is_index_type()) { @@ -31,7 +31,7 @@ struct dispatch_create_indices { } template std::unique_ptr operator()(column_view const&, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref) requires(!is_index_type()) { @@ -42,7 +42,7 @@ struct dispatch_create_indices { std::unique_ptr make_dictionary_column(column_view const& keys_column, column_view const& indices_column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(!keys_column.has_nulls(), "keys column must not have nulls", std::invalid_argument); @@ -114,7 +114,7 @@ struct make_signed_fn { std::unique_ptr make_dictionary_column(std::unique_ptr keys, std::unique_ptr indices, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(!keys->has_nulls(), "keys column must not have nulls", std::invalid_argument); diff --git a/cpp/src/dictionary/encode.cu b/cpp/src/dictionary/encode.cu index a43d18c080a4..7f43eb337370 100644 --- a/cpp/src/dictionary/encode.cu +++ b/cpp/src/dictionary/encode.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 */ @@ -23,7 +23,6 @@ #include #include -#include #include #include #include @@ -32,6 +31,7 @@ #include #include #include +#include #include #include #include @@ -55,7 +55,7 @@ struct encode_fn { std::unique_ptr encode(column_view const& input, data_type indices_type, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(is_signed(indices_type) && is_index_type(indices_type), @@ -89,8 +89,8 @@ std::unique_ptr encode(column_view const& input, auto const empty_key = cuco::empty_key{cudf::detail::CUDF_SIZE_TYPE_SENTINEL}; auto probe = encode_probe_t{row_hash.device_hasher(has_nulls)}; auto allocator = rmm::mr::polymorphic_allocator{}; - auto set = cuco::static_set{ - input.size(), 0.5, empty_key, d_equal, probe, {}, {}, allocator, stream.value()}; + auto set = + cuco::static_set{input.size(), 0.5, empty_key, d_equal, probe, {}, {}, allocator, stream.get()}; auto set_ref = set.ref(cuco::insert_and_find); using set_ref_t = decltype(set_ref); @@ -105,7 +105,7 @@ std::unique_ptr encode(column_view const& input, encode_fn{set_ref, *d_input}); auto keys_indices = rmm::device_uvector(input.size(), stream); - auto keys_end = set.retrieve_all(keys_indices.begin(), stream.value()); + auto keys_end = set.retrieve_all(keys_indices.begin(), stream.get()); keys_indices.resize(cuda::std::distance(keys_indices.begin(), keys_end), stream); // sort the keys_indices so we can use lower-bound on them @@ -154,7 +154,7 @@ data_type get_indices_type_for_size(size_type keys_size) std::unique_ptr encode(column_view const& input_column, data_type indices_type, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/dictionary/match_keys.cu b/cpp/src/dictionary/match_keys.cu index 3637d231d41d..269cb403998c 100644 --- a/cpp/src/dictionary/match_keys.cu +++ b/cpp/src/dictionary/match_keys.cu @@ -19,13 +19,13 @@ #include #include -#include #include #include #include #include #include +#include #include #include @@ -39,7 +39,7 @@ namespace { struct unique_keys_dispatch_fn { template std::unique_ptr operator()(cudf::column_view const& all_keys, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(cudf::is_dictionary_key()) { @@ -58,15 +58,15 @@ struct unique_keys_dispatch_fn { auto probe = probe_t{row_hash.device_hasher(has_nulls)}; auto allocator = rmm::mr::polymorphic_allocator{}; auto set = cuco::static_set{ - all_keys.size(), 0.5, empty_key, d_equal, probe, {}, {}, allocator, stream.value()}; + all_keys.size(), 0.5, empty_key, d_equal, probe, {}, {}, allocator, stream.get()}; // use a static_set to find the unique elements of all_keys auto const iter = cuda::counting_iterator{0}; - set.insert_async(iter, iter + all_keys.size(), stream.value()); + set.insert_async(iter, iter + all_keys.size(), stream.get()); // retrieve the indices of all the unique keys auto keys_indices = rmm::device_uvector(all_keys.size(), stream); - auto keys_end = set.retrieve_all(keys_indices.begin(), stream.value()); + auto keys_end = set.retrieve_all(keys_indices.begin(), stream.get()); keys_indices.resize(cuda::std::distance(keys_indices.begin(), keys_end), stream); // gather the unique keys using the keys_indices @@ -81,7 +81,7 @@ struct unique_keys_dispatch_fn { template std::unique_ptr operator()(cudf::column_view const&, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref) requires(not cudf::is_dictionary_key()) { @@ -92,7 +92,7 @@ struct unique_keys_dispatch_fn { std::vector> match_dictionaries( std::span input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(not input.empty(), "expect at least one dictionary", std::invalid_argument); @@ -114,7 +114,7 @@ std::vector> match_dictionaries( } std::pair>, std::vector> match_dictionaries( - std::vector tables, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + std::vector tables, cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // Make a copy of all the column views from each table_view std::vector> updated_columns; @@ -164,7 +164,7 @@ std::pair>, std::vector> match_d std::vector> match_dictionaries( std::span input, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/dictionary/remove_keys.cu b/cpp/src/dictionary/remove_keys.cu index 91912dd6b253..7ae4826b6a3d 100644 --- a/cpp/src/dictionary/remove_keys.cu +++ b/cpp/src/dictionary/remove_keys.cu @@ -20,10 +20,10 @@ #include #include -#include #include #include +#include #include #include #include @@ -51,7 +51,7 @@ namespace { template std::unique_ptr remove_keys_fn(dictionary_column_view const& dictionary_column, KeysKeeper keys_to_keep_fn, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const keys_view = dictionary_column.keys(); @@ -150,7 +150,7 @@ std::unique_ptr remove_keys_fn(dictionary_column_view const& dictionary_ std::unique_ptr remove_keys(dictionary_column_view const& dictionary_column, column_view const& keys_to_remove, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(!keys_to_remove.has_nulls(), "keys_to_remove must not have nulls"); @@ -168,7 +168,7 @@ std::unique_ptr remove_keys(dictionary_column_view const& dictionary_col } std::unique_ptr remove_unused_keys(dictionary_column_view const& dictionary_column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // locate the keys to remove @@ -200,7 +200,7 @@ std::unique_ptr remove_unused_keys(dictionary_column_view const& diction std::unique_ptr remove_keys(dictionary_column_view const& dictionary_column, column_view const& keys_to_remove, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -208,7 +208,7 @@ std::unique_ptr remove_keys(dictionary_column_view const& dictionary_col } std::unique_ptr remove_unused_keys(dictionary_column_view const& dictionary_column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/dictionary/replace.cu b/cpp/src/dictionary/replace.cu index f722a64afe2f..77a5b17a381d 100644 --- a/cpp/src/dictionary/replace.cu +++ b/cpp/src/dictionary/replace.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 */ @@ -17,7 +17,7 @@ #include #include -#include +#include namespace cudf { namespace dictionary { @@ -43,7 +43,7 @@ namespace { template std::unique_ptr replace_indices(column_view const& input, ReplacementIter replacement_iter, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const input_view = column_device_view::create(input, stream); @@ -66,11 +66,11 @@ std::unique_ptr replace_indices(column_view const& input, /** * @copydoc cudf::dictionary::detail::replace_nulls(cudf::column_view const&,cudf::column_view - * const& rmm::cuda_stream_view, rmm::device_async_resource_ref) + * const& cuda::stream_ref, rmm::device_async_resource_ref) */ std::unique_ptr replace_nulls(dictionary_column_view const& input, dictionary_column_view 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.parent()); } @@ -101,11 +101,11 @@ std::unique_ptr replace_nulls(dictionary_column_view const& input, /** * @copydoc cudf::dictionary::detail::replace_nulls(cudf::column_view const&,cudf::scalar - * const&, rmm::cuda_stream_view, rmm::device_async_resource_ref) + * const&, cuda::stream_ref, rmm::device_async_resource_ref) */ std::unique_ptr replace_nulls(dictionary_column_view const& input, 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.parent()); } diff --git a/cpp/src/dictionary/search.cu b/cpp/src/dictionary/search.cu index 9bd4b6002df5..d0e868572907 100644 --- a/cpp/src/dictionary/search.cu +++ b/cpp/src/dictionary/search.cu @@ -15,9 +15,8 @@ #include #include -#include - #include +#include namespace cudf { namespace dictionary { @@ -36,7 +35,7 @@ struct find_index_fn { template std::unique_ptr operator()(dictionary_column_view const& input, scalar const& key, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const requires(not std::is_same_v and not std::is_same_v and not std::is_same_v) @@ -60,10 +59,10 @@ struct find_index_fn { auto find_fn = [find_key] __device__(auto const& k) { return k == find_key.value(); }; auto tmp_size = std::size_t{0}; CUDF_CUDA_TRY(cub::DeviceFind::FindIf( - nullptr, tmp_size, keys, result->data(), find_fn, num_keys, stream.value())); + nullptr, tmp_size, keys, result->data(), find_fn, num_keys, stream.get())); auto tmp = rmm::device_buffer(tmp_size, stream); CUDF_CUDA_TRY(cub::DeviceFind::FindIf( - tmp.data(), tmp_size, keys, result->data(), find_fn, num_keys, stream.value())); + tmp.data(), tmp_size, keys, result->data(), find_fn, num_keys, stream.get())); if (result->value(stream) == num_keys) { result->set_valid_async(false, stream); } return result; } @@ -71,7 +70,7 @@ struct find_index_fn { template std::unique_ptr operator()(dictionary_column_view const&, scalar const&, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref) const requires(std::is_same_v or std::is_same_v or std::is_same_v) @@ -85,7 +84,7 @@ struct find_index_fn { std::unique_ptr get_index(dictionary_column_view const& dictionary, scalar const& key, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (dictionary.is_empty()) { @@ -101,7 +100,7 @@ std::unique_ptr get_index(dictionary_column_view const& dictionary, std::unique_ptr get_index(dictionary_column_view const& dictionary, scalar const& key, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/dictionary/set_keys.cu b/cpp/src/dictionary/set_keys.cu index 79f22c9d463a..04475f2bbcdb 100644 --- a/cpp/src/dictionary/set_keys.cu +++ b/cpp/src/dictionary/set_keys.cu @@ -22,11 +22,11 @@ #include #include -#include #include #include #include +#include #include #include #include @@ -69,7 +69,7 @@ struct set_keys_dispatch_fn { template std::unique_ptr operator()(cudf::dictionary_column_view const& input, cudf::column_view const& new_keys, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(cudf::is_dictionary_key()) { @@ -124,7 +124,7 @@ struct set_keys_dispatch_fn { template std::unique_ptr operator()(cudf::dictionary_column_view const&, cudf::column_view const&, - rmm::cuda_stream_view, + cuda::stream_ref, rmm::device_async_resource_ref) requires(not cudf::is_dictionary_key()) { @@ -135,7 +135,7 @@ struct set_keys_dispatch_fn { std::unique_ptr set_keys(dictionary_column_view const& input, column_view const& new_keys, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(!new_keys.has_nulls(), "keys parameter must not have nulls", std::invalid_argument); @@ -153,7 +153,7 @@ std::unique_ptr set_keys(dictionary_column_view const& input, std::unique_ptr set_keys(dictionary_column_view const& dictionary_column, column_view 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/filling/calendrical_month_sequence.cu b/cpp/src/filling/calendrical_month_sequence.cu index 893645aea051..b020b4817701 100644 --- a/cpp/src/filling/calendrical_month_sequence.cu +++ b/cpp/src/filling/calendrical_month_sequence.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -12,15 +12,16 @@ #include #include -#include #include +#include + namespace cudf { namespace detail { std::unique_ptr calendrical_month_sequence(size_type size, scalar const& init, size_type months, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return type_dispatcher( @@ -31,7 +32,7 @@ std::unique_ptr calendrical_month_sequence(size_type size, std::unique_ptr calendrical_month_sequence(size_type size, scalar const& init, size_type months, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/filling/fill.cu b/cpp/src/filling/fill.cu index cd15422a43b1..fb80d3e67113 100644 --- a/cpp/src/filling/fill.cu +++ b/cpp/src/filling/fill.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 */ @@ -25,9 +25,8 @@ #include #include -#include - #include +#include #include @@ -37,7 +36,7 @@ void in_place_fill(cudf::mutable_column_view& destination, cudf::size_type begin, cudf::size_type end, cudf::scalar const& value, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { using ScalarType = cudf::scalar_type_t; auto p_scalar = static_cast(&value); @@ -57,7 +56,7 @@ struct in_place_fill_range_dispatch { template std::enable_if_t() && not cudf::is_fixed_point(), void> operator()( - cudf::size_type begin, cudf::size_type end, rmm::cuda_stream_view stream) + cudf::size_type begin, cudf::size_type end, cuda::stream_ref stream) { in_place_fill(destination, begin, end, value, stream); } @@ -65,7 +64,7 @@ struct in_place_fill_range_dispatch { template std::enable_if_t(), void> operator()(cudf::size_type begin, cudf::size_type end, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto unscaled = static_cast const&>(value).value(stream); using RepType = typename T::rep; @@ -97,7 +96,7 @@ struct out_of_place_fill_range_dispatch { CUDF_ENABLE_IF(cudf::is_rep_layout_compatible() or cudf::is_fixed_point())> std::unique_ptr operator()(cudf::size_type begin, cudf::size_type end, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(cudf::have_same_types(input, value), "Data type mismatch.", cudf::data_type_error); @@ -124,7 +123,7 @@ template <> std::unique_ptr out_of_place_fill_range_dispatch::operator()( cudf::size_type begin, cudf::size_type end, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(cudf::have_same_types(input, value), "Data type mismatch.", cudf::data_type_error); @@ -138,7 +137,7 @@ template <> std::unique_ptr out_of_place_fill_range_dispatch::operator()( cudf::size_type begin, cudf::size_type end, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (input.is_empty()) return std::make_unique(input, stream, mr); @@ -189,7 +188,7 @@ void fill_in_place(mutable_column_view& destination, size_type begin, size_type end, scalar const& value, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_EXPECTS(cudf::is_fixed_width(destination.type()), "In-place fill does not support variable-sized types."); @@ -212,7 +211,7 @@ std::unique_ptr fill(column_view const& input, size_type begin, size_type end, scalar const& value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS((begin >= 0) && (end <= input.size()) && (begin <= end), "Range is out of bounds."); @@ -227,7 +226,7 @@ void fill_in_place(mutable_column_view& destination, size_type begin, size_type end, scalar const& value, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { CUDF_FUNC_RANGE(); return detail::fill_in_place(destination, begin, end, value, stream); @@ -237,7 +236,7 @@ std::unique_ptr fill(column_view const& input, size_type begin, size_type end, scalar const& value, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/filling/repeat.cu b/cpp/src/filling/repeat.cu index 2ba8b2539af3..8d328ae581f9 100644 --- a/cpp/src/filling/repeat.cu +++ b/cpp/src/filling/repeat.cu @@ -19,12 +19,12 @@ #include #include -#include #include #include #include #include +#include #include #include #include @@ -38,7 +38,7 @@ struct count_accessor { cudf::scalar const* p_scalar = nullptr; template - cudf::size_type operator()(rmm::cuda_stream_view stream) + cudf::size_type operator()(cuda::stream_ref stream) requires(std::is_integral_v) { using ScalarType = cudf::scalar_type_t; @@ -52,7 +52,7 @@ struct count_accessor { } template - cudf::size_type operator()(rmm::cuda_stream_view) + cudf::size_type operator()(cuda::stream_ref) requires(not std::is_integral_v) { CUDF_FAIL("count value should be a integral type."); @@ -63,7 +63,7 @@ struct count_checker { cudf::column_view const& count; template - void operator()(rmm::cuda_stream_view stream) + void operator()(cuda::stream_ref stream) requires(std::is_integral_v) { // static_cast is necessary due to bool @@ -82,7 +82,7 @@ struct count_checker { } template - void operator()(rmm::cuda_stream_view) + void operator()(cuda::stream_ref) requires(not std::is_integral_v) { CUDF_FAIL("count value type should be integral."); @@ -95,7 +95,7 @@ namespace cudf { namespace detail { std::unique_ptr
repeat(table_view const& input_table, column_view const& count, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(input_table.num_rows() == count.size(), "in and count must have equal size"); @@ -126,7 +126,7 @@ std::unique_ptr
repeat(table_view const& input_table, std::unique_ptr
repeat(table_view const& input_table, size_type count, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if ((input_table.num_rows() == 0) || (count == 0)) { return cudf::empty_like(input_table); } @@ -149,7 +149,7 @@ std::unique_ptr
repeat(table_view const& input_table, std::unique_ptr
repeat(table_view const& input_table, column_view const& count, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -158,7 +158,7 @@ std::unique_ptr
repeat(table_view const& input_table, std::unique_ptr
repeat(table_view const& input_table, 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/filling/sequence.cu b/cpp/src/filling/sequence.cu index b8fa91bf6b18..67f66b908470 100644 --- a/cpp/src/filling/sequence.cu +++ b/cpp/src/filling/sequence.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 */ @@ -17,9 +17,9 @@ #include #include -#include #include +#include #include #include @@ -56,7 +56,7 @@ struct sequence_functor { std::unique_ptr operator()(size_type size, scalar const& init, scalar const& step, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(cudf::is_numeric() and not cudf::is_boolean()) { @@ -82,7 +82,7 @@ struct sequence_functor { template std::unique_ptr operator()(size_type size, scalar const& init, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) requires(cudf::is_numeric() and not cudf::is_boolean()) { @@ -116,7 +116,7 @@ struct sequence_functor { std::unique_ptr sequence(size_type size, scalar const& init, scalar const& step, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(cudf::have_same_types(init, step), @@ -133,7 +133,7 @@ std::unique_ptr sequence(size_type size, std::unique_ptr sequence(size_type size, scalar const& init, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(size >= 0, "size must be >= 0", std::invalid_argument); @@ -148,7 +148,7 @@ std::unique_ptr sequence(size_type size, std::unique_ptr sequence(size_type size, scalar const& init, scalar const& step, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -157,7 +157,7 @@ std::unique_ptr sequence(size_type size, std::unique_ptr sequence(size_type size, scalar const& init, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/partitioning/partitioning.cu b/cpp/src/partitioning/partitioning.cu index d9bf0d99864f..62a3586e2ebc 100644 --- a/cpp/src/partitioning/partitioning.cu +++ b/cpp/src/partitioning/partitioning.cu @@ -20,7 +20,6 @@ #include #include -#include #include #include @@ -29,6 +28,7 @@ #include #include #include +#include #include #include @@ -346,7 +346,7 @@ void copy_block_partitions_impl(InputIter const input, size_type const* block_partition_sizes, size_type const* scanned_block_partition_sizes, size_type grid_size, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // We need 3 chunks of shared memory: // 1. BLOCK_SIZE * ROWS_PER_THREAD elements of size_type for copying to output @@ -355,7 +355,7 @@ void copy_block_partitions_impl(InputIter const input, int const smem = OPTIMIZED_BLOCK_SIZE * OPTIMIZED_ROWS_PER_THREAD * sizeof(*output) + (num_partitions + 1) * sizeof(size_type) * 2; - copy_block_partitions<<>>( + copy_block_partitions<<>>( input, output, num_rows, @@ -374,7 +374,7 @@ rmm::device_uvector compute_gather_map(size_type num_rows, size_type const* block_partition_sizes, size_type const* scanned_block_partition_sizes, size_type grid_size, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto sequence = cuda::counting_iterator{0}; rmm::device_uvector gather_map(num_rows, stream); @@ -410,7 +410,7 @@ struct copy_block_partitions_dispatcher { size_type const* block_partition_sizes, size_type const* scanned_block_partition_sizes, size_type grid_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { rmm::device_buffer output(input.size() * sizeof(DataType), stream, mr); @@ -438,7 +438,7 @@ struct copy_block_partitions_dispatcher { size_type const* block_partition_sizes, size_type const* scanned_block_partition_sizes, size_type grid_size, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // Use move_to_output_buffer to create an equivalent gather map @@ -470,7 +470,7 @@ std::pair, std::vector> hash_partition_table_g size_type num_rows, size_type num_partitions, Hasher hasher, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(num_partitions < std::numeric_limits::max(), @@ -518,7 +518,7 @@ std::pair, std::vector> hash_partition_table_g lower_level, upper_level, num_rows, - stream.value()); + stream.get()); rmm::device_buffer temp_storage(temp_storage_bytes, stream); cub::DeviceHistogram::HistogramEven(temp_storage.data(), temp_storage_bytes, @@ -528,7 +528,7 @@ std::pair, std::vector> hash_partition_table_g lower_level, upper_level, num_rows, - stream.value()); + stream.get()); } // Exclusive scan on histogram to get partition offsets. @@ -556,7 +556,7 @@ std::pair, std::vector> hash_partition_table_g // Scatter input rows into partitioned output auto output = detail::scatter(input, scatter_map, input, stream, mr); - stream.synchronize(); // Pinned async D2H copy must finish before returning host vec + stream.wait(); // Pinned async D2H copy must finish before returning host vec // Convert pinned host_vector to std::vector for the return type auto partition_offsets = std::vector(pinned_offsets.begin(), pinned_offsets.end()); @@ -571,7 +571,7 @@ std::pair, std::vector> hash_partition_table( table_view const& table_to_hash, size_type num_partitions, uint32_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const num_rows = table_to_hash.num_rows(); @@ -634,14 +634,14 @@ std::pair, std::vector> hash_partition_table( compute_row_partition_numbers<<>>(hasher, - num_rows, - num_partitions, - partitioner_type(num_partitions), - row_partition_numbers.data(), - row_partition_offset.data(), - block_partition_sizes.data(), - global_partition_sizes.data()); + stream.get()>>>(hasher, + num_rows, + num_partitions, + partitioner_type(num_partitions), + row_partition_numbers.data(), + row_partition_offset.data(), + block_partition_sizes.data(), + global_partition_sizes.data()); CUDF_CUDA_TRY(cudaGetLastError()); } else { // Determines how the mapping between hash value and partition number is @@ -655,14 +655,14 @@ std::pair, std::vector> hash_partition_table( compute_row_partition_numbers<<>>(hasher, - num_rows, - num_partitions, - partitioner_type(num_partitions), - row_partition_numbers.data(), - row_partition_offset.data(), - block_partition_sizes.data(), - global_partition_sizes.data()); + stream.get()>>>(hasher, + num_rows, + num_partitions, + partitioner_type(num_partitions), + row_partition_numbers.data(), + row_partition_offset.data(), + block_partition_sizes.data(), + global_partition_sizes.data()); CUDF_CUDA_TRY(cudaGetLastError()); } @@ -724,7 +724,7 @@ std::pair, std::vector> hash_partition_table( input, gather_map.begin(), output_cols, detail::gather_bitmask_op::DONT_CHECK, stream, mr); } - stream.synchronize(); // Async D2H copy must finish before returning host vec + stream.wait(); // Async D2H copy must finish before returning host vec return std::pair{std::make_unique
(std::move(output_cols), num_rows), std::move(partition_offsets)}; } else { @@ -735,14 +735,14 @@ std::pair, std::vector> hash_partition_table( compute_row_output_locations<<>>( + stream.get()>>>( row_output_locations, num_rows, num_partitions, scanned_block_partition_sizes_ptr); CUDF_CUDA_TRY(cudaGetLastError()); // Use the resulting scatter map to materialize the output auto output = detail::scatter(input, row_partition_numbers, input, stream, mr); - stream.synchronize(); // Async D2H copy must finish before returning host vec + stream.wait(); // Async D2H copy must finish before returning host vec return std::pair{std::move(output), std::move(partition_offsets)}; } } @@ -772,7 +772,7 @@ struct dispatch_map_type { table_view const& t, column_view const& partition_map, size_type num_partitions, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) const requires(is_index_type()) { @@ -790,7 +790,7 @@ struct dispatch_map_type { lower_level, upper_level, partition_map.size(), - stream.value()); + stream.get()); rmm::device_buffer temp_storage(temp_storage_bytes, stream); @@ -802,7 +802,7 @@ struct dispatch_map_type { lower_level, upper_level, partition_map.size(), - stream.value()); + stream.get()); // `histogram` was created with an extra entry at the end such that an // exclusive scan will put the total number of rows at the end @@ -877,7 +877,7 @@ std::pair, std::vector> hash_partition( table_view const& table_to_hash, int num_partitions, uint32_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // Return empty result if there are no partitions or nothing to hash @@ -904,7 +904,7 @@ std::pair, std::vector> partition( table_view const& t, column_view const& partition_map, size_type num_partitions, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(t.num_rows() == partition_map.size(), @@ -926,7 +926,7 @@ std::pair, std::vector> hash_partition( int num_partitions, hash_id hash_function, uint32_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS( @@ -951,7 +951,7 @@ std::pair, std::vector> hash_partition( int num_partitions, hash_id hash_function, uint32_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -965,7 +965,7 @@ std::pair, std::vector> hash_partition( int num_partitions, hash_id hash_function, uint32_t seed, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); @@ -977,7 +977,7 @@ std::pair, std::vector> partition( table_view const& t, column_view const& partition_map, size_type num_partitions, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/partitioning/round_robin.cu b/cpp/src/partitioning/round_robin.cu index 5a6d3080fa0e..92559da84276 100644 --- a/cpp/src/partitioning/round_robin.cu +++ b/cpp/src/partitioning/round_robin.cu @@ -19,12 +19,12 @@ #include #include -#include #include #include #include #include +#include #include #include #include @@ -72,7 +72,7 @@ std::pair, std::vector> degenerate cudf::table_view const& input, cudf::size_type num_partitions, cudf::size_type start_partition, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto nrows = input.num_rows(); @@ -152,7 +152,7 @@ std::pair, std::vector> round_robin_part table_view const& input, cudf::size_type num_partitions, cudf::size_type start_partition, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto nrows = input.num_rows(); @@ -269,7 +269,7 @@ std::pair, std::vector> round_robi table_view const& input, cudf::size_type num_partitions, cudf::size_type start_partition, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/scalar/scalar.cpp b/cpp/src/scalar/scalar.cpp index cd437c360496..97945228c6f7 100644 --- a/cpp/src/scalar/scalar.cpp +++ b/cpp/src/scalar/scalar.cpp @@ -12,34 +12,35 @@ #include #include -#include #include +#include + #include namespace cudf { scalar::scalar(data_type type, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : _type(type), _is_valid(is_valid, stream, mr) { } -scalar::scalar(scalar const& other, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) +scalar::scalar(scalar const& other, cuda::stream_ref stream, rmm::device_async_resource_ref mr) : _type(other.type()), _is_valid(other._is_valid, stream, mr) { } data_type scalar::type() const noexcept { return _type; } -void scalar::set_valid_async(bool is_valid, rmm::cuda_stream_view stream) +void scalar::set_valid_async(bool is_valid, cuda::stream_ref stream) { _is_valid.set_value_async(is_valid, stream); } -bool scalar::is_valid(rmm::cuda_stream_view stream) const { return _is_valid.value(stream); } +bool scalar::is_valid(cuda::stream_ref stream) const { return _is_valid.value(stream); } bool* scalar::validity_data() { return _is_valid.data(); } @@ -47,7 +48,7 @@ bool const* scalar::validity_data() const { return _is_valid.data(); } string_scalar::string_scalar(std::string_view string, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar(data_type(type_id::STRING), is_valid, stream, mr), _data(string.data(), string.size(), stream, mr) @@ -59,7 +60,7 @@ string_scalar::string_scalar(std::string_view string, } string_scalar::string_scalar(string_scalar const& other, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar(other, stream, mr), _data(other._data, stream, mr) { @@ -67,7 +68,7 @@ string_scalar::string_scalar(string_scalar const& other, string_scalar::string_scalar(rmm::device_scalar& data, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : string_scalar(data.value(stream), is_valid, stream, mr) { @@ -75,7 +76,7 @@ string_scalar::string_scalar(rmm::device_scalar& data, string_scalar::string_scalar(value_type const& source, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar(data_type(type_id::STRING), is_valid, stream, mr), _data(source.data(), source.size_bytes(), stream, mr) @@ -84,13 +85,13 @@ string_scalar::string_scalar(value_type const& source, string_scalar::string_scalar(rmm::device_buffer&& data, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar(data_type(type_id::STRING), is_valid, stream, mr), _data(std::move(data)) { } -string_scalar::value_type string_scalar::value(rmm::cuda_stream_view stream) const +string_scalar::value_type string_scalar::value(cuda::stream_ref stream) const { return value_type{data(), size()}; } @@ -99,7 +100,7 @@ size_type string_scalar::size() const { return _data.size(); } char const* string_scalar::data() const { return static_cast(_data.data()); } -std::string string_scalar::to_string(rmm::cuda_stream_view stream) const +std::string string_scalar::to_string(cuda::stream_ref stream) const { std::string result(size(), '\0'); detail::cuda_memcpy(host_span{result.data(), result.size()}, @@ -112,7 +113,7 @@ template fixed_point_scalar::fixed_point_scalar(rep_type value, numeric::scale_type scale, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar{data_type{type_to_id(), static_cast(scale)}, is_valid, stream, mr}, _data{value, stream, mr} @@ -122,7 +123,7 @@ fixed_point_scalar::fixed_point_scalar(rep_type value, template fixed_point_scalar::fixed_point_scalar(rep_type value, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar{data_type{type_to_id(), 0}, is_valid, stream, mr}, _data{value, stream, mr} { @@ -131,7 +132,7 @@ fixed_point_scalar::fixed_point_scalar(rep_type value, template fixed_point_scalar::fixed_point_scalar(T value, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar{data_type{type_to_id(), value.scale()}, is_valid, stream, mr}, _data{value.value(), stream, mr} @@ -142,7 +143,7 @@ template fixed_point_scalar::fixed_point_scalar(rmm::device_scalar&& data, numeric::scale_type scale, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar{data_type{type_to_id(), scale}, is_valid, stream, mr}, _data{std::move(data)} { @@ -150,21 +151,20 @@ fixed_point_scalar::fixed_point_scalar(rmm::device_scalar&& data, template fixed_point_scalar::fixed_point_scalar(fixed_point_scalar const& other, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar{other, stream, mr}, _data(other._data, stream, mr) { } template -typename fixed_point_scalar::rep_type fixed_point_scalar::value( - rmm::cuda_stream_view stream) const +typename fixed_point_scalar::rep_type fixed_point_scalar::value(cuda::stream_ref stream) const { return _data.value(stream); } template -T fixed_point_scalar::fixed_point_value(rmm::cuda_stream_view stream) const +T fixed_point_scalar::fixed_point_value(cuda::stream_ref stream) const { return value_type{ numeric::scaled_integer{_data.value(stream), numeric::scale_type{type().scale()}}}; @@ -199,7 +199,7 @@ namespace CUDF_HIDDEN detail { template fixed_width_scalar::fixed_width_scalar(T value, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar(data_type(type_to_id()), is_valid, stream, mr), _data(value, stream, mr) { @@ -208,7 +208,7 @@ fixed_width_scalar::fixed_width_scalar(T value, template fixed_width_scalar::fixed_width_scalar(rmm::device_scalar&& data, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar(data_type(type_to_id()), is_valid, stream, mr), _data{std::move(data)} { @@ -216,21 +216,21 @@ fixed_width_scalar::fixed_width_scalar(rmm::device_scalar&& data, template fixed_width_scalar::fixed_width_scalar(fixed_width_scalar const& other, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar{other, stream, mr}, _data(other._data, stream, mr) { } template -void fixed_width_scalar::set_value(T value, rmm::cuda_stream_view stream) +void fixed_width_scalar::set_value(T value, cuda::stream_ref stream) { _data.set_value_async(value, stream); this->set_valid_async(true, stream); } template -T fixed_width_scalar::value(rmm::cuda_stream_view stream) const +T fixed_width_scalar::value(cuda::stream_ref stream) const { return _data.value(stream); } @@ -283,7 +283,7 @@ template class fixed_width_scalar; template numeric_scalar::numeric_scalar(T value, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : detail::fixed_width_scalar(value, is_valid, stream, mr) { @@ -292,7 +292,7 @@ numeric_scalar::numeric_scalar(T value, template numeric_scalar::numeric_scalar(rmm::device_scalar&& data, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : detail::fixed_width_scalar(std::forward>(data), is_valid, stream, mr) { @@ -300,7 +300,7 @@ numeric_scalar::numeric_scalar(rmm::device_scalar&& data, template numeric_scalar::numeric_scalar(numeric_scalar const& other, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : detail::fixed_width_scalar{other, stream, mr} { @@ -330,7 +330,7 @@ template class numeric_scalar; template chrono_scalar::chrono_scalar(T value, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : detail::fixed_width_scalar(value, is_valid, stream, mr) { @@ -339,7 +339,7 @@ chrono_scalar::chrono_scalar(T value, template chrono_scalar::chrono_scalar(rmm::device_scalar&& data, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : detail::fixed_width_scalar(std::forward>(data), is_valid, stream, mr) { @@ -347,7 +347,7 @@ chrono_scalar::chrono_scalar(rmm::device_scalar&& data, template chrono_scalar::chrono_scalar(chrono_scalar const& other, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : detail::fixed_width_scalar{other, stream, mr} { @@ -375,7 +375,7 @@ template class chrono_scalar; template duration_scalar::duration_scalar(rep_type value, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : chrono_scalar(T{value}, is_valid, stream, mr) { @@ -383,14 +383,14 @@ duration_scalar::duration_scalar(rep_type value, template duration_scalar::duration_scalar(duration_scalar const& other, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : chrono_scalar{other, stream, mr} { } template -typename duration_scalar::rep_type duration_scalar::count(rmm::cuda_stream_view stream) +typename duration_scalar::rep_type duration_scalar::count(cuda::stream_ref stream) { return this->value(stream).count(); } @@ -411,7 +411,7 @@ template class duration_scalar; template typename timestamp_scalar::rep_type timestamp_scalar::ticks_since_epoch( - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { return this->value(stream).time_since_epoch().count(); } @@ -434,7 +434,7 @@ template template timestamp_scalar::timestamp_scalar(D const& value, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : chrono_scalar(T{typename T::duration{value}}, is_valid, stream, mr) { @@ -442,7 +442,7 @@ timestamp_scalar::timestamp_scalar(D const& value, template timestamp_scalar::timestamp_scalar(timestamp_scalar const& other, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : chrono_scalar{other, stream, mr} { @@ -450,7 +450,7 @@ timestamp_scalar::timestamp_scalar(timestamp_scalar const& other, #define TS_CTOR(TimestampType, DurationType) \ template timestamp_scalar::timestamp_scalar( \ - DurationType const&, bool, rmm::cuda_stream_view, rmm::device_async_resource_ref); + DurationType const&, bool, cuda::stream_ref, rmm::device_async_resource_ref); /** * @brief These are the valid combinations of duration types to timestamp types. @@ -478,7 +478,7 @@ TS_CTOR(timestamp_ns, int64_t) list_scalar::list_scalar(cudf::column_view const& data, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar(data_type(type_id::LIST), is_valid, stream, mr), _data(data, stream, mr) { @@ -486,14 +486,14 @@ list_scalar::list_scalar(cudf::column_view const& data, list_scalar::list_scalar(cudf::column&& data, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar(data_type(type_id::LIST), is_valid, stream, mr), _data(std::move(data)) { } list_scalar::list_scalar(list_scalar const& other, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar{other, stream, mr}, _data(other._data, stream, mr) { @@ -502,7 +502,7 @@ list_scalar::list_scalar(list_scalar const& other, column_view list_scalar::view() const { return _data.view(); } struct_scalar::struct_scalar(struct_scalar const& other, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar{other, stream, mr}, _data(other._data, stream, mr) { @@ -510,7 +510,7 @@ struct_scalar::struct_scalar(struct_scalar const& other, struct_scalar::struct_scalar(table_view const& data, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar(data_type(type_id::STRUCT), is_valid, stream, mr), _data{init_data(table{data, stream, mr}, is_valid, stream, mr)} @@ -520,7 +520,7 @@ struct_scalar::struct_scalar(table_view const& data, struct_scalar::struct_scalar(std::span data, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar(data_type(type_id::STRUCT), is_valid, stream, mr), _data{ @@ -534,7 +534,7 @@ struct_scalar::struct_scalar(std::span data, struct_scalar::struct_scalar(table&& data, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) : scalar(data_type(type_id::STRUCT), is_valid, stream, mr), _data{init_data(std::move(data), is_valid, stream, mr)} @@ -554,7 +554,7 @@ void struct_scalar::assert_valid_size() table struct_scalar::init_data(table&& data, bool is_valid, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { if (is_valid) { return std::move(data); } diff --git a/cpp/src/scalar/scalar_factories.cpp b/cpp/src/scalar/scalar_factories.cpp index 08b9f8ec4459..8216723b7afe 100644 --- a/cpp/src/scalar/scalar_factories.cpp +++ b/cpp/src/scalar/scalar_factories.cpp @@ -10,7 +10,7 @@ #include #include -#include +#include namespace cudf { namespace { @@ -19,7 +19,7 @@ struct scalar_construction_helper { template () and not is_fixed_point()>* = nullptr> - std::unique_ptr operator()(rmm::cuda_stream_view stream, + std::unique_ptr operator()(cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { using Type = device_storage_type_t; @@ -28,7 +28,7 @@ struct scalar_construction_helper { } template ()>* = nullptr> - std::unique_ptr operator()(rmm::cuda_stream_view stream, + std::unique_ptr operator()(cuda::stream_ref stream, rmm::device_async_resource_ref mr) const { using Type = device_storage_type_t; @@ -47,7 +47,7 @@ struct scalar_construction_helper { // Allocate storage for a single numeric element std::unique_ptr make_numeric_scalar(data_type type, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(is_numeric(type), "Invalid, non-numeric type."); @@ -57,7 +57,7 @@ std::unique_ptr make_numeric_scalar(data_type type, // Allocate storage for a single timestamp element std::unique_ptr make_timestamp_scalar(data_type type, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(is_timestamp(type), "Invalid, non-timestamp type."); @@ -67,7 +67,7 @@ std::unique_ptr make_timestamp_scalar(data_type type, // Allocate storage for a single duration element std::unique_ptr make_duration_scalar(data_type type, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(is_duration(type), "Invalid, non-duration type."); @@ -77,7 +77,7 @@ std::unique_ptr make_duration_scalar(data_type type, // Allocate storage for a single fixed width element std::unique_ptr make_fixed_width_scalar(data_type type, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(is_fixed_width(type), "Invalid, non-fixed-width type."); @@ -86,21 +86,21 @@ std::unique_ptr make_fixed_width_scalar(data_type type, } std::unique_ptr make_list_scalar(column_view elements, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return std::make_unique(elements, true, stream, mr); } std::unique_ptr make_struct_scalar(table_view const& data, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return std::make_unique(data, true, stream, mr); } std::unique_ptr make_struct_scalar(std::span data, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return std::make_unique(data, true, stream, mr); @@ -111,14 +111,14 @@ struct default_scalar_functor { data_type type; template ()>* = nullptr> - std::unique_ptr operator()(rmm::cuda_stream_view stream, + std::unique_ptr operator()(cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return make_fixed_width_scalar(data_type(type_to_id()), stream, mr); } template ()>* = nullptr> - std::unique_ptr operator()(rmm::cuda_stream_view stream, + std::unique_ptr operator()(cuda::stream_ref stream, rmm::device_async_resource_ref mr) { auto const scale_ = numeric::scale_type{type.scale()}; @@ -130,28 +130,28 @@ struct default_scalar_functor { template <> std::unique_ptr default_scalar_functor::operator()( - rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return std::unique_ptr(new string_scalar("", false, stream, mr)); } template <> std::unique_ptr default_scalar_functor::operator()( - rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FAIL("dictionary type not supported"); } template <> std::unique_ptr default_scalar_functor::operator()( - rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FAIL("list_view type not supported"); } template <> std::unique_ptr default_scalar_functor::operator()( - rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FAIL("struct_view type not supported"); } @@ -159,14 +159,14 @@ std::unique_ptr default_scalar_functor::operator()( } // namespace std::unique_ptr make_default_constructed_scalar(data_type type, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { return type_dispatcher(type, default_scalar_functor{type}, stream, mr); } std::unique_ptr make_empty_scalar_like(column_view const& column, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { std::unique_ptr result; diff --git a/cpp/src/search/contains_column.cu b/cpp/src/search/contains_column.cu index 639b31b9b47b..1a9c6d46ace7 100644 --- a/cpp/src/search/contains_column.cu +++ b/cpp/src/search/contains_column.cu @@ -12,7 +12,7 @@ #include #include -#include +#include namespace cudf { namespace detail { @@ -21,7 +21,7 @@ namespace { std::unique_ptr contains_dictionary(column_view const& haystack_in, column_view const& needles_in, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { dictionary_column_view const haystack(haystack_in); @@ -52,7 +52,7 @@ std::unique_ptr contains_dictionary(column_view const& haystack_in, std::unique_ptr contains(column_view const& haystack, column_view const& needles, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { // Dictionary columns require key normalization; all other types share the type-erased path. @@ -73,7 +73,7 @@ std::unique_ptr contains(column_view const& haystack, std::unique_ptr contains(column_view const& haystack, column_view const& needles, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_FUNC_RANGE(); diff --git a/cpp/src/search/contains_scalar.cu b/cpp/src/search/contains_scalar.cu index 3f70f0260f69..c6f329e5e1f8 100644 --- a/cpp/src/search/contains_scalar.cu +++ b/cpp/src/search/contains_scalar.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 */ @@ -20,9 +20,9 @@ #include #include -#include #include +#include #include namespace cudf { @@ -53,7 +53,7 @@ struct contains_scalar_dispatch { template std::enable_if_t(), bool> operator()(column_view const& haystack, scalar const& needle, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { CUDF_EXPECTS(cudf::have_same_types(haystack, needle), "Scalar and column types must match", @@ -83,7 +83,7 @@ struct contains_scalar_dispatch { template std::enable_if_t(), bool> operator()(column_view const& haystack, scalar const& needle, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { CUDF_EXPECTS(cudf::have_same_types(haystack, needle), "Scalar and column types must match", @@ -137,7 +137,7 @@ struct contains_scalar_dispatch { template <> bool contains_scalar_dispatch::operator()(column_view const& haystack, scalar const& needle, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { auto const dict_col = cudf::dictionary_column_view(haystack); // first, find the needle in the dictionary's key set @@ -153,7 +153,7 @@ bool contains_scalar_dispatch::operator()(column_view const& } // namespace -bool contains(column_view const& haystack, scalar const& needle, rmm::cuda_stream_view stream) +bool contains(column_view const& haystack, scalar const& needle, cuda::stream_ref stream) { if (haystack.is_empty()) { return false; } if (not needle.is_valid(stream)) { return haystack.has_nulls(); } @@ -164,7 +164,7 @@ bool contains(column_view const& haystack, scalar const& needle, rmm::cuda_strea } // namespace detail -bool contains(column_view const& haystack, scalar const& needle, rmm::cuda_stream_view stream) +bool contains(column_view const& haystack, scalar const& needle, cuda::stream_ref stream) { CUDF_FUNC_RANGE(); return detail::contains(haystack, needle, stream); diff --git a/cpp/src/search/contains_table.cu b/cpp/src/search/contains_table.cu index c3ff93adf4a3..189199ef1cf6 100644 --- a/cpp/src/search/contains_table.cu +++ b/cpp/src/search/contains_table.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -13,10 +13,10 @@ #include #include -#include #include #include +#include #include @@ -26,7 +26,7 @@ rmm::device_uvector contains(table_view const& haystack, table_view const& needles, null_equality compare_nulls, nan_equality compare_nans, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, rmm::device_async_resource_ref mr) { CUDF_EXPECTS(cudf::have_same_types(haystack, needles), "Column types mismatch"); diff --git a/cpp/src/search/contains_table_impl.cu b/cpp/src/search/contains_table_impl.cu index cf8caff1107d..a4733f0ec52a 100644 --- a/cpp/src/search/contains_table_impl.cu +++ b/cpp/src/search/contains_table_impl.cu @@ -20,7 +20,7 @@ namespace cudf::detail { * @return A pair of pointer to the output bitmask and the buffer containing the bitmask */ std::pair build_row_bitmask(table_view const& input, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const nullable_columns = get_nullable_columns(input); CUDF_EXPECTS(nullable_columns.size() > 0, diff --git a/cpp/src/search/contains_table_impl.cuh b/cpp/src/search/contains_table_impl.cuh index 792b9aa000f1..3d8f1f5df9c5 100644 --- a/cpp/src/search/contains_table_impl.cuh +++ b/cpp/src/search/contains_table_impl.cuh @@ -15,13 +15,13 @@ #include #include -#include #include #include #include #include #include +#include namespace cudf::detail { @@ -94,7 +94,7 @@ struct comparator_adapter { * @return A pair of pointer to the output bitmask and the buffer containing the bitmask */ std::pair build_row_bitmask(table_view const& input, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); /** * @brief Helper function to perform the contains operation using a hash set @@ -121,7 +121,7 @@ void perform_contains(table_view const& haystack, Comparator const& d_equal, ProbingScheme const& probing_scheme, rmm::device_uvector& contained, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto const haystack_iter = cudf::detail::make_counting_transform_iterator( size_type{0}, cuda::proclaim_return_type([] __device__(auto idx) { @@ -141,7 +141,7 @@ void perform_contains(table_view const& haystack, {}, {}, rmm::mr::polymorphic_allocator{}, - stream.value()}; + stream.get()}; if (haystack_has_nulls && compare_nulls == null_equality::UNEQUAL) { auto const bitmask_buffer_and_ptr = build_row_bitmask(haystack, stream); @@ -155,9 +155,9 @@ void perform_contains(table_view const& haystack, haystack_iter + haystack.num_rows(), cuda::counting_iterator{0}, // stencil row_is_valid{row_bitmask_ptr}, - stream.value()); + stream.get()); } else { - set.insert_async(haystack_iter, haystack_iter + haystack.num_rows(), stream.value()); + set.insert_async(haystack_iter, haystack_iter + haystack.num_rows(), stream.get()); } if (needles_has_nulls && compare_nulls == null_equality::UNEQUAL) { @@ -168,10 +168,10 @@ void perform_contains(table_view const& haystack, cuda::counting_iterator{0}, // stencil row_is_valid{row_bitmask_ptr}, contained.begin(), - stream.value()); + stream.get()); } else { set.contains_async( - needles_iter, needles_iter + needles.num_rows(), contained.begin(), stream.value()); + needles_iter, needles_iter + needles.num_rows(), contained.begin(), stream.get()); } } @@ -207,7 +207,7 @@ void dispatch_nan_comparator(table_view const& haystack, cudf::detail::row::equality::two_table_comparator two_table_equal, Hasher const& d_hasher, rmm::device_uvector& contained, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { // Distinguish probing scheme CG sizes between nested and flat types for better performance auto const probing_scheme = [&]() { diff --git a/cpp/src/search/search_ordered.cu b/cpp/src/search/search_ordered.cu index 1f4cfa38c019..b579d9f40975 100644 --- a/cpp/src/search/search_ordered.cu +++ b/cpp/src/search/search_ordered.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 */ @@ -14,9 +14,9 @@ #include #include -#include #include +#include #include namespace cudf { @@ -28,7 +28,7 @@ std::unique_ptr search_ordered(table_view const& haystack, bool find_first, 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_EXPECTS( @@ -45,8 +45,7 @@ std::unique_ptr search_ordered(table_view const& haystack, // Handle empty inputs if (haystack.num_rows() == 0) { - CUDF_CUDA_TRY( - cudaMemsetAsync(out_it, 0, needles.num_rows() * sizeof(size_type), stream.value())); + CUDF_CUDA_TRY(cudaMemsetAsync(out_it, 0, needles.num_rows() * sizeof(size_type), stream.get())); return result; } @@ -111,7 +110,7 @@ std::unique_ptr lower_bound(table_view const& haystack, table_view const& needles, 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) { return search_ordered(haystack, needles, true, column_order, null_precedence, stream, mr); @@ -121,7 +120,7 @@ std::unique_ptr upper_bound(table_view const& haystack, table_view const& needles, 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) { return search_ordered(haystack, needles, false, column_order, null_precedence, stream, mr); @@ -135,7 +134,7 @@ std::unique_ptr lower_bound(table_view const& haystack, table_view const& needles, 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(); @@ -146,7 +145,7 @@ std::unique_ptr upper_bound(table_view const& haystack, table_view const& needles, 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/tests/column/column_device_view_test.cu b/cpp/tests/column/column_device_view_test.cu index 96ebff66e0d4..b7608c3fac73 100644 --- a/cpp/tests/column/column_device_view_test.cu +++ b/cpp/tests/column/column_device_view_test.cu @@ -16,9 +16,9 @@ #include #include -#include #include +#include #include struct ColumnDeviceViewTest : public cudf::test::BaseFixture {}; @@ -26,7 +26,7 @@ struct ColumnDeviceViewTest : public cudf::test::BaseFixture {}; TEST_F(ColumnDeviceViewTest, Sample) { using T = int32_t; - rmm::cuda_stream_view stream{cudf::get_default_stream()}; + cuda::stream_ref stream{cudf::get_default_stream()}; cudf::test::fixed_width_column_wrapper input({1, 2, 3, 4, 5, 6}); auto output = cudf::allocate_like(input); auto input_device_view = cudf::column_device_view::create(input, stream); @@ -44,7 +44,7 @@ TEST_F(ColumnDeviceViewTest, Sample) TEST_F(ColumnDeviceViewTest, MismatchingType) { using T = int32_t; - rmm::cuda_stream_view stream{cudf::get_default_stream()}; + cuda::stream_ref stream{cudf::get_default_stream()}; cudf::test::fixed_width_column_wrapper input({1, 2, 3, 4, 5, 6}); auto output = cudf::allocate_like(input); auto input_device_view = cudf::column_device_view::create(input, stream); diff --git a/cpp/tests/column/column_test.cpp b/cpp/tests/column/column_test.cpp index 2b3da507b858..50fe302d0b1b 100644 --- a/cpp/tests/column/column_test.cpp +++ b/cpp/tests/column/column_test.cpp @@ -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 */ @@ -35,16 +35,16 @@ template struct TypedColumnTest : public cudf::test::BaseFixture { cudf::data_type type() { return cudf::data_type{cudf::type_to_id()}; } - TypedColumnTest(rmm::cuda_stream_view stream = cudf::get_default_stream()) + TypedColumnTest(cuda::stream_ref stream = cudf::get_default_stream()) : data{_num_elements * sizeof(T), stream}, mask{cudf::bitmask_allocation_size_bytes(_num_elements), stream} { std::vector h_data(std::max(data.size(), mask.size())); std::iota(h_data.begin(), h_data.end(), 0); CUDF_CUDA_TRY( - cudaMemcpyAsync(data.data(), h_data.data(), data.size(), cudaMemcpyDefault, stream.value())); + cudaMemcpyAsync(data.data(), h_data.data(), data.size(), cudaMemcpyDefault, stream.get())); CUDF_CUDA_TRY( - cudaMemcpyAsync(mask.data(), h_data.data(), mask.size(), cudaMemcpyDefault, stream.value())); + cudaMemcpyAsync(mask.data(), h_data.data(), mask.size(), cudaMemcpyDefault, stream.get())); } cudf::size_type num_elements() { return _num_elements; } diff --git a/cpp/tests/copying/concatenate_tests.cpp b/cpp/tests/copying/concatenate_tests.cpp index 8875add31856..4ad0afe255fa 100644 --- a/cpp/tests/copying/concatenate_tests.cpp +++ b/cpp/tests/copying/concatenate_tests.cpp @@ -45,7 +45,7 @@ template struct TypedColumnTest : public cudf::test::BaseFixture { cudf::data_type type() { return cudf::data_type{cudf::type_to_id()}; } - TypedColumnTest(rmm::cuda_stream_view stream = cudf::get_default_stream()) + TypedColumnTest(cuda::stream_ref stream = cudf::get_default_stream()) : data{_num_elements * sizeof(T), stream}, mask{cudf::bitmask_allocation_size_bytes(_num_elements), stream} { @@ -56,11 +56,11 @@ struct TypedColumnTest : public cudf::test::BaseFixture { std::vector h_mask(mask.size()); std::iota(h_mask.begin(), h_mask.end(), char{0}); CUDF_CUDA_TRY( - cudaMemcpyAsync(typed_data, h_data.data(), data.size(), cudaMemcpyDefault, stream.value())); + cudaMemcpyAsync(typed_data, h_data.data(), data.size(), cudaMemcpyDefault, stream.get())); CUDF_CUDA_TRY( - cudaMemcpyAsync(typed_mask, h_mask.data(), mask.size(), cudaMemcpyDefault, stream.value())); + cudaMemcpyAsync(typed_mask, h_mask.data(), mask.size(), cudaMemcpyDefault, stream.get())); _null_count = cudf::null_count(static_cast(mask.data()), 0, _num_elements); - stream.synchronize(); + stream.wait(); } [[nodiscard]] cudf::size_type num_elements() const { return _num_elements; } diff --git a/cpp/tests/copying/shift_tests.cpp b/cpp/tests/copying/shift_tests.cpp index 1a865c3e3d21..ae4a12e4f12b 100644 --- a/cpp/tests/copying/shift_tests.cpp +++ b/cpp/tests/copying/shift_tests.cpp @@ -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 */ @@ -14,7 +14,7 @@ #include #include -#include +#include #include #include @@ -23,7 +23,7 @@ using TestTypes = cudf::test::Types; template > std::unique_ptr make_scalar( - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()) { auto s = new ScalarType(cudf::test::make_type_param_scalar(0), false, stream, mr); @@ -33,7 +33,7 @@ std::unique_ptr make_scalar( template > std::unique_ptr make_scalar( T value, - rmm::cuda_stream_view stream = cudf::get_default_stream(), + cuda::stream_ref stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()) { auto s = new ScalarType(value, true, stream, mr); diff --git a/cpp/tests/streams/column_view_test.cpp b/cpp/tests/streams/column_view_test.cpp index 8db9fa659b75..0c0541311b9c 100644 --- a/cpp/tests/streams/column_view_test.cpp +++ b/cpp/tests/streams/column_view_test.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -20,16 +20,16 @@ template struct TypedColumnTest : public cudf::test::BaseFixture { cudf::data_type type() { return cudf::data_type{cudf::type_to_id()}; } - TypedColumnTest(rmm::cuda_stream_view stream = cudf::test::get_default_stream()) + TypedColumnTest(cuda::stream_ref stream = cudf::test::get_default_stream()) : data{_num_elements * sizeof(T), stream}, mask{cudf::bitmask_allocation_size_bytes(_num_elements), stream} { std::vector h_data(std::max(data.size(), mask.size())); std::iota(h_data.begin(), h_data.end(), 0); CUDF_CUDA_TRY( - cudaMemcpyAsync(data.data(), h_data.data(), data.size(), cudaMemcpyDefault, stream.value())); + cudaMemcpyAsync(data.data(), h_data.data(), data.size(), cudaMemcpyDefault, stream.get())); CUDF_CUDA_TRY( - cudaMemcpyAsync(mask.data(), h_data.data(), mask.size(), cudaMemcpyDefault, stream.value())); + cudaMemcpyAsync(mask.data(), h_data.data(), mask.size(), cudaMemcpyDefault, stream.get())); } cudf::size_type num_elements() { return _num_elements; } diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu index 48d664272d16..046ebcbb917b 100644 --- a/cpp/tests/utilities/column_utilities.cu +++ b/cpp/tests/utilities/column_utilities.cu @@ -30,6 +30,7 @@ #include #include #include +#include #include #include #include @@ -53,7 +54,7 @@ namespace test { namespace { std::unique_ptr generate_all_row_indices(size_type num_rows, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { auto indices = cudf::make_fixed_width_column( @@ -91,7 +92,7 @@ std::unique_ptr generate_all_row_indices(size_type num_rows, std::unique_ptr generate_child_row_indices(lists_column_view const& c, column_view const& row_indices, bool check_exact_equality, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { // if we are checking for exact equality, we should be checking for "unsanitized" data that may @@ -242,7 +243,7 @@ struct column_property_comparator { cudf::column_view const& lhs_row_indices, cudf::column_view const& rhs_row_indices, debug_output_level verbosity, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { bool result = true; @@ -276,7 +277,7 @@ struct column_property_comparator { cudf::column_view const& lhs_row_indices, cudf::column_view const& rhs_row_indices, debug_output_level verbosity, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) requires(!std::is_same_v && !std::is_same_v) { @@ -289,7 +290,7 @@ struct column_property_comparator { cudf::column_view const& lhs_row_indices, cudf::column_view const& rhs_row_indices, debug_output_level verbosity, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) requires(std::is_same_v) { @@ -330,7 +331,7 @@ struct column_property_comparator { cudf::column_view const& lhs_row_indices, cudf::column_view const& rhs_row_indices, debug_output_level verbosity, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) requires(std::is_same_v) { @@ -477,7 +478,7 @@ std::string stringify_column_differences(cudf::device_span difference column_view const& rhs_row_indices, debug_output_level verbosity, int depth, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { CUDF_EXPECTS(not differences.empty(), "Shouldn't enter this function if `differences` is empty"); @@ -527,7 +528,7 @@ struct column_comparator_impl { debug_output_level verbosity, size_type fp_ulps, int depth, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { auto d_lhs_row_indices = @@ -607,7 +608,7 @@ struct column_comparator_impl { debug_output_level verbosity, size_type fp_ulps, int depth, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { lists_column_view lhs_l(lhs); @@ -752,7 +753,7 @@ struct column_comparator_impl { debug_output_level verbosity, size_type fp_ulps, int depth, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { structs_column_view l_scv(lhs); @@ -789,7 +790,7 @@ struct column_comparator { debug_output_level verbosity, size_type fp_ulps, int depth, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { // compare properties @@ -812,9 +813,7 @@ struct column_comparator { } }; -void check_non_empty_nulls(column_view const& lhs, - column_view const& rhs, - rmm::cuda_stream_view stream) +void check_non_empty_nulls(column_view const& lhs, column_view const& rhs, cuda::stream_ref stream) { auto check_column_nulls = [stream](column_view const& col, char const* col_name) { if (cudf::detail::has_nonempty_nulls(col, stream)) { @@ -835,7 +834,7 @@ namespace detail { bool expect_column_properties_equal(column_view const& lhs, column_view const& rhs, debug_output_level verbosity, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { check_non_empty_nulls(lhs, rhs, stream); @@ -858,7 +857,7 @@ bool expect_column_properties_equal(column_view const& lhs, bool expect_column_properties_equivalent(column_view const& lhs, column_view const& rhs, debug_output_level verbosity, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { check_non_empty_nulls(lhs, rhs, stream); @@ -881,7 +880,7 @@ bool expect_column_properties_equivalent(column_view const& lhs, bool expect_columns_equal(cudf::column_view const& lhs, cudf::column_view const& rhs, debug_output_level verbosity, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { check_non_empty_nulls(lhs, rhs, stream); @@ -907,7 +906,7 @@ bool expect_columns_equivalent(cudf::column_view const& lhs, cudf::column_view const& rhs, debug_output_level verbosity, size_type fp_ulps, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { check_non_empty_nulls(lhs, rhs, stream); @@ -932,7 +931,7 @@ bool expect_columns_equivalent(cudf::column_view const& lhs, void expect_equal_buffers(void const* lhs, void const* rhs, std::size_t size_bytes, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { if (size_bytes > 0) { @@ -961,7 +960,7 @@ void expect_column_empty(cudf::column_view const& col) * @copydoc cudf::test::bitmask_to_host */ std::vector bitmask_to_host(cudf::column_view const& c, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { if (c.nullable()) { @@ -1000,7 +999,7 @@ bool validate_host_masks(std::vector const& expected_mask, template ()>*> std::pair, std::vector> to_host(column_view c, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, cudf::memory_resources mr) { using namespace numeric; @@ -1017,11 +1016,11 @@ std::pair, std::vector> to_host(column_view } template std::pair, std::vector> to_host( - column_view c, rmm::cuda_stream_view stream, cudf::memory_resources mr); + column_view c, cuda::stream_ref stream, cudf::memory_resources mr); template std::pair, std::vector> to_host( - column_view c, rmm::cuda_stream_view stream, cudf::memory_resources mr); + column_view c, cuda::stream_ref stream, cudf::memory_resources mr); template std::pair, std::vector> to_host( - column_view c, rmm::cuda_stream_view stream, cudf::memory_resources mr); + column_view c, cuda::stream_ref stream, cudf::memory_resources mr); namespace { struct strings_to_host_fn { @@ -1029,7 +1028,7 @@ struct strings_to_host_fn { void operator()(thrust::host_vector& host_data, char const* chars, cudf::column_view const& offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) requires(std::is_same_v || std::is_same_v) { auto const h_offsets = cudf::detail::make_std_vector( @@ -1046,7 +1045,7 @@ struct strings_to_host_fn { void operator()(thrust::host_vector&, char const*, cudf::column_view const&, - rmm::cuda_stream_view) + cuda::stream_ref) requires(!std::is_same_v && !std::is_same_v) { CUDF_FAIL("invalid offsets type"); @@ -1056,7 +1055,7 @@ struct strings_to_host_fn { template <> std::pair, std::vector> to_host( - column_view c, rmm::cuda_stream_view stream, cudf::memory_resources mr) + column_view c, cuda::stream_ref stream, cudf::memory_resources mr) { thrust::host_vector host_data(c.size()); if (c.size() > c.null_count()) {