Use cuda::stream_ref in remaining core libcudf APIs

cpp/include/cudf/ast/detail/expression_evaluator.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -16,10 +16,9 @@
 #include <cudf/utilities/span.hpp>
 #include <cudf/utilities/traits.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
-
 #include <cuda/std/type_traits>
 #include <cuda/std/utility>
+#include <cuda/stream>
 
 namespace cudf::ast::detail {
 

cpp/include/cudf/ast/detail/expression_parser.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -106,7 +106,7 @@ class expression_parser {
                     cudf::table_view const& left,
                     std::optional<std::reference_wrapper<cudf::table_view const>> right,
                     bool has_nulls,
-                    rmm::cuda_stream_view stream,
+                    cuda::stream_ref stream,
                     rmm::device_async_resource_ref mr);
 
   /**
@@ -121,7 +121,7 @@ class expression_parser {
   expression_parser(expression const& expr,
                     cudf::table_view const& table,
                     bool has_nulls,
-                    rmm::cuda_stream_view stream,
+                    cuda::stream_ref stream,
                     rmm::device_async_resource_ref mr);
 
   /**
@@ -240,7 +240,7 @@ class expression_parser {
     alignment = std::max(alignment, static_cast<cudf::size_type>(alignof(T)));
   }
 
-  void move_to_device(rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr);
+  void move_to_device(cuda::stream_ref stream, rmm::device_async_resource_ref mr);
 
   /**
    * @brief Helper function for recursive traversal of expressions.

cpp/include/cudf/ast/expressions.hpp

@@ -13,6 +13,8 @@
 #include <cudf/types.hpp>
 #include <cudf/utilities/error.hpp>
 
+#include <cuda/stream>
+
 #include <cstdint>
 #include <functional>
 #include <initializer_list>
@@ -101,7 +103,7 @@ struct [[nodiscard]] expression {
    * @param stream CUDA stream used for device memory operations and kernel launches
    * @return `true` if the expression may evaluate to null, otherwise false
    */
-  [[nodiscard]] bool may_evaluate_null(table_view const& left, rmm::cuda_stream_view stream) const
+  [[nodiscard]] bool may_evaluate_null(table_view const& left, cuda::stream_ref stream) const
   {
     return may_evaluate_null(left, left, stream);
   }
@@ -116,7 +118,7 @@ struct [[nodiscard]] expression {
    */
   [[nodiscard]] virtual bool may_evaluate_null(table_view const& left,
                                                table_view const& right,
-                                               rmm::cuda_stream_view stream) const = 0;
+                                               cuda::stream_ref stream) const = 0;
 
   virtual ~expression() {}
 };
@@ -384,7 +386,7 @@ class literal : public expression {
 
   [[nodiscard]] bool may_evaluate_null(table_view const& left,
                                        table_view const& right,
-                                       rmm::cuda_stream_view stream) const override
+                                       cuda::stream_ref stream) const override
   {
     return !is_valid(stream);
   }
@@ -395,7 +397,7 @@ class literal : public expression {
    * @param stream CUDA stream used for device memory operations and kernel launches
    * @return true if the underlying scalar is valid
    */
-  [[nodiscard]] bool is_valid(rmm::cuda_stream_view stream) const
+  [[nodiscard]] bool is_valid(cuda::stream_ref stream) const
   {
     if (auto* s = std::get_if<ast_scalar>(&scalar)) {
       return s->scalar.get().is_valid(stream);
@@ -477,7 +479,7 @@ class column_reference : public expression {
 
   [[nodiscard]] bool may_evaluate_null(table_view const& left,
                                        table_view const& right,
-                                       rmm::cuda_stream_view stream) const override
+                                       cuda::stream_ref stream) const override
   {
     return (table_source == table_reference::LEFT ? left : right).column(column_index).has_nulls();
   }
@@ -552,7 +554,7 @@ class operation : public expression {
 
   [[nodiscard]] bool may_evaluate_null(table_view const& left,
                                        table_view const& right,
-                                       rmm::cuda_stream_view stream) const override;
+                                       cuda::stream_ref stream) const override;
 
   /**
    * @copydoc expression::accept
@@ -592,7 +594,7 @@ class predicate : public expression {
 
   [[nodiscard]] bool may_evaluate_null(table_view const& left,
                                        table_view const& right,
-                                       rmm::cuda_stream_view stream) const override;
+                                       cuda::stream_ref stream) const override;
 
   /**
    * @copydoc expression::accept
@@ -645,7 +647,7 @@ class column_name_reference : public expression {
 
   [[nodiscard]] bool may_evaluate_null(table_view const& left,
                                        table_view const& right,
-                                       rmm::cuda_stream_view stream) const override
+                                       cuda::stream_ref stream) const override
   {
     return true;
   }

cpp/include/cudf/datetime.hpp

@@ -10,7 +10,7 @@
 #include <cudf/utilities/export.hpp>
 #include <cudf/utilities/memory_resource.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
+#include <cuda/stream>
 
 #include <memory>
 
@@ -57,7 +57,7 @@ enum class datetime_component : uint8_t {
 std::unique_ptr<cudf::column> extract_datetime_component(
   cudf::column_view const& column,
   datetime_component component,
-  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
+  cuda::stream_ref stream           = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
 /** @} */  // end of group
@@ -79,7 +79,7 @@ std::unique_ptr<cudf::column> extract_datetime_component(
  */
 std::unique_ptr<cudf::column> last_day_of_month(
   cudf::column_view const& column,
-  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
+  cuda::stream_ref stream           = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
 /**
@@ -95,7 +95,7 @@ std::unique_ptr<cudf::column> last_day_of_month(
  */
 std::unique_ptr<cudf::column> day_of_year(
   cudf::column_view const& column,
-  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
+  cuda::stream_ref stream           = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
 /**
@@ -132,7 +132,7 @@ std::unique_ptr<cudf::column> day_of_year(
 std::unique_ptr<cudf::column> add_calendrical_months(
   cudf::column_view const& timestamps,
   cudf::column_view const& months,
-  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
+  cuda::stream_ref stream           = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
 /**
@@ -169,7 +169,7 @@ std::unique_ptr<cudf::column> add_calendrical_months(
 std::unique_ptr<cudf::column> add_calendrical_months(
   cudf::column_view const& timestamps,
   cudf::scalar const& months,
-  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
+  cuda::stream_ref stream           = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
 /**
@@ -188,7 +188,7 @@ std::unique_ptr<cudf::column> add_calendrical_months(
  */
 std::unique_ptr<cudf::column> is_leap_year(
   cudf::column_view const& column,
-  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
+  cuda::stream_ref stream           = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
 /**
@@ -207,7 +207,7 @@ std::unique_ptr<cudf::column> is_leap_year(
  */
 std::unique_ptr<cudf::column> days_in_month(
   cudf::column_view const& column,
-  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
+  cuda::stream_ref stream           = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
 /**
@@ -227,7 +227,7 @@ std::unique_ptr<cudf::column> days_in_month(
  */
 std::unique_ptr<cudf::column> extract_quarter(
   cudf::column_view const& column,
-  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
+  cuda::stream_ref stream           = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
 /**
@@ -258,7 +258,7 @@ enum class rounding_frequency : int32_t {
 std::unique_ptr<cudf::column> ceil_datetimes(
   cudf::column_view const& column,
   rounding_frequency freq,
-  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
+  cuda::stream_ref stream           = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
 /**
@@ -275,7 +275,7 @@ std::unique_ptr<cudf::column> ceil_datetimes(
 std::unique_ptr<cudf::column> floor_datetimes(
   cudf::column_view const& column,
   rounding_frequency freq,
-  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
+  cuda::stream_ref stream           = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
 /**
@@ -292,7 +292,7 @@ std::unique_ptr<cudf::column> floor_datetimes(
 std::unique_ptr<cudf::column> round_datetimes(
   cudf::column_view const& column,
   rounding_frequency freq,
-  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
+  cuda::stream_ref stream           = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
 /** @} */  // end of group

cpp/include/cudf/detail/aggregation/aggregation.hpp

@@ -1382,7 +1382,7 @@ bool is_valid_aggregation(data_type source, aggregation::Kind k);
  */
 void initialize_with_identity(mutable_table_view const& table,
                               host_span<cudf::aggregation::Kind const> aggs,
-                              rmm::cuda_stream_view stream);
+                              cuda::stream_ref stream);
 
 }  // namespace detail
 }  // namespace cudf

cpp/include/cudf/detail/algorithms/copy_if.cuh

@@ -8,7 +8,6 @@
 #include <cudf/utilities/error.hpp>
 #include <cudf/utilities/memory_resource.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_buffer.hpp>
 #include <rmm/exec_policy.hpp>
 
@@ -49,7 +48,7 @@ OutputIterator copy_if(InputIterator begin,
                        StencilIterator stencil,
                        OutputIterator result,
                        Predicate predicate,
-                       rmm::cuda_stream_view stream)
+                       cuda::stream_ref stream)
 {
   auto const num_items = cuda::std::distance(begin, end);
 
@@ -65,7 +64,7 @@ OutputIterator copy_if(InputIterator begin,
                                              num_selected.data(),
                                              num_items,
                                              predicate,
-                                             stream.value()));
+                                             stream.get()));
 
   auto d_temp_storage =
     rmm::device_buffer(temp_storage_bytes, stream, cudf::get_current_device_resource_ref());
@@ -78,7 +77,7 @@ OutputIterator copy_if(InputIterator begin,
                                              num_selected.data(),
                                              num_items,
                                              predicate,
-                                             stream.value()));
+                                             stream.get()));
 
   return result + num_selected.value(stream);
 }
@@ -106,7 +105,7 @@ OutputIterator copy_if(InputIterator begin,
                        InputIterator end,
                        OutputIterator output,
                        Predicate predicate,
-                       rmm::cuda_stream_view stream)
+                       cuda::stream_ref stream)
 {
   auto const num_items = cuda::std::distance(begin, end);
 
@@ -123,7 +122,7 @@ OutputIterator copy_if(InputIterator begin,
                                       num_selected.data(),
                                       num_items,
                                       predicate,
-                                      stream.value()));
+                                      stream.get()));
 
   // Allocate temporary storage
   rmm::device_buffer d_temp_storage(
@@ -137,7 +136,7 @@ OutputIterator copy_if(InputIterator begin,
                                       num_selected.data(),
                                       num_items,
                                       predicate,
-                                      stream.value()));
+                                      stream.get()));
 
   // Copy number of selected elements back to host via pinned memory
   return output + num_selected.value(stream);
@@ -155,18 +154,18 @@ void copy_if_async(InputIterator begin,
                    InputIterator end,
                    OutputIterator output,
                    Predicate predicate,
-                   rmm::cuda_stream_view stream)
+                   cuda::stream_ref stream)
 {
   auto const num_items = cuda::std::distance(begin, end);
 
   auto tmp_bytes = std::size_t{0};
   auto no_out    = cuda::make_discard_iterator<int>();
   CUDF_CUDA_TRY(cub::DeviceSelect::If(
-    nullptr, tmp_bytes, begin, output, no_out, num_items, predicate, stream.value()));
+    nullptr, tmp_bytes, begin, output, no_out, num_items, predicate, stream.get()));
 
   auto tmp_stg = rmm::device_buffer(tmp_bytes, stream, cudf::get_current_device_resource_ref());
   CUDF_CUDA_TRY(cub::DeviceSelect::If(
-    tmp_stg.data(), tmp_bytes, begin, output, no_out, num_items, predicate, stream.value()));
+    tmp_stg.data(), tmp_bytes, begin, output, no_out, num_items, predicate, stream.get()));
 }
 
 /**
@@ -185,18 +184,18 @@ void copy_if_async(InputIterator begin,
                    StencilIterator stencil,
                    OutputIterator result,
                    Predicate predicate,
-                   rmm::cuda_stream_view stream)
+                   cuda::stream_ref stream)
 {
   auto const num_items = cuda::std::distance(begin, end);
 
   auto tmp_bytes = std::size_t{0};
   auto no_out    = cuda::make_discard_iterator<int>();
   CUDF_CUDA_TRY(cub::DeviceSelect::FlaggedIf(
-    nullptr, tmp_bytes, begin, stencil, result, no_out, num_items, predicate, stream.value()));
+    nullptr, tmp_bytes, begin, stencil, result, no_out, num_items, predicate, stream.get()));
 
   auto tmp = rmm::device_buffer(tmp_bytes, stream, cudf::get_current_device_resource_ref());
   CUDF_CUDA_TRY(cub::DeviceSelect::FlaggedIf(
-    tmp.data(), tmp_bytes, begin, stencil, result, no_out, num_items, predicate, stream.value()));
+    tmp.data(), tmp_bytes, begin, stencil, result, no_out, num_items, predicate, stream.get()));
 }
 
 }  // namespace cudf::detail