diff --git a/cpp/docs/DEVELOPER_GUIDE.md b/cpp/docs/DEVELOPER_GUIDE.md
index ba074b0e88..7ae90b844d 100644
--- a/cpp/docs/DEVELOPER_GUIDE.md
+++ b/cpp/docs/DEVELOPER_GUIDE.md
@@ -134,7 +134,7 @@ Similar to a `rmm::device_vector`, allocates a contiguous set of elements in dev
 key differences:
 - As an optimization, elements are uninitialized and no synchronization occurs at construction.
 This limits the types `T` to trivially copyable types.
-- All operations are stream ordered (i.e., they accept a `cuda_stream_view` specifying the stream
+- All operations are stream ordered (i.e., they accept a `cuda::stream_ref` specifying the stream
 on which the operation is performed).
 
 ## Namespaces
diff --git a/cpp/include/cuopt/mathematical_optimization/cpu_pdlp_warm_start_data.hpp b/cpp/include/cuopt/mathematical_optimization/cpu_pdlp_warm_start_data.hpp
index 1a76da0fa3..a1ea73858a 100644
--- a/cpp/include/cuopt/mathematical_optimization/cpu_pdlp_warm_start_data.hpp
+++ b/cpp/include/cuopt/mathematical_optimization/cpu_pdlp_warm_start_data.hpp
@@ -7,6 +7,7 @@
 
 #pragma once
 
+#include <cuda/stream>
 #include <cuopt/export.hpp>
 #include <cuopt/mathematical_optimization/pdlp/pdlp_warm_start_data.hpp>
 
@@ -114,12 +115,12 @@ struct pdlp_warm_start_data_t;
 // Convert GPU → CPU warmstart (D2H copy)
 template <typename i_t, typename f_t>
 cpu_pdlp_warm_start_data_t<i_t, f_t> convert_to_cpu_warmstart(
-  const pdlp_warm_start_data_t<i_t, f_t>& gpu_data, rmm::cuda_stream_view stream);
+  const pdlp_warm_start_data_t<i_t, f_t>& gpu_data, cuda::stream_ref stream);
 
 // Convert CPU → GPU warmstart (H2D copy)
 template <typename i_t, typename f_t>
 pdlp_warm_start_data_t<i_t, f_t> convert_to_gpu_warmstart(
-  const cpu_pdlp_warm_start_data_t<i_t, f_t>& cpu_data, rmm::cuda_stream_view stream);
+  const cpu_pdlp_warm_start_data_t<i_t, f_t>& cpu_data, cuda::stream_ref stream);
 
 }  // namespace CUOPT_EXPORT mathematical_optimization
 }  // namespace cuopt
diff --git a/cpp/include/cuopt/mathematical_optimization/mip/solver_settings.hpp b/cpp/include/cuopt/mathematical_optimization/mip/solver_settings.hpp
index f04c676442..7ed45f1f9b 100644
--- a/cpp/include/cuopt/mathematical_optimization/mip/solver_settings.hpp
+++ b/cpp/include/cuopt/mathematical_optimization/mip/solver_settings.hpp
@@ -20,7 +20,6 @@
 #include <cuopt/mathematical_optimization/utilities/internals.hpp>
 
 #include <cuda/stream>
-
 #include <raft/core/device_span.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -92,7 +91,7 @@ class mip_solver_settings_t {
    */
   void add_initial_solution(const f_t* initial_solution,
                             i_t size,
-                            rmm::cuda_stream_view stream = cuda::stream_ref{
+                            cuda::stream_ref stream = cuda::stream_ref{
                               cudaStream_t{cudaStreamDefault}});
 
   /**
diff --git a/cpp/include/cuopt/mathematical_optimization/mip/solver_solution.hpp b/cpp/include/cuopt/mathematical_optimization/mip/solver_solution.hpp
index 1ad58b9e10..510382d2f7 100644
--- a/cpp/include/cuopt/mathematical_optimization/mip/solver_solution.hpp
+++ b/cpp/include/cuopt/mathematical_optimization/mip/solver_solution.hpp
@@ -13,7 +13,7 @@
 #include <cuopt/mathematical_optimization/mip/solver_stats.hpp>
 #include <cuopt/mathematical_optimization/utilities/internals.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
+#include <cuda/stream>
 #include <rmm/device_uvector.hpp>
 
 #include <raft/core/handle.hpp>
@@ -52,8 +52,8 @@ class mip_solution_t : public base_solution_t {
 
   mip_solution_t(mip_termination_status_t termination_status,
                  solver_stats_t<i_t, f_t> stats,
-                 rmm::cuda_stream_view stream_view);
-  mip_solution_t(const cuopt::logic_error& error_status, rmm::cuda_stream_view stream_view);
+                 cuda::stream_ref stream_view);
+  mip_solution_t(const cuopt::logic_error& error_status, cuda::stream_ref stream_view);
 
   bool is_mip() const override { return true; }
   const rmm::device_uvector<f_t>& get_solution() const;
@@ -76,7 +76,7 @@ class mip_solution_t : public base_solution_t {
   i_t get_num_simplex_iterations() const;
   const std::vector<std::string>& get_variable_names() const;
   const std::vector<rmm::device_uvector<f_t>>& get_solution_pool() const;
-  void write_to_sol_file(std::string_view filename, rmm::cuda_stream_view stream_view) const;
+  void write_to_sol_file(std::string_view filename, cuda::stream_ref stream_view) const;
   void log_detailed_summary() const;
   void log_summary() const;
 
diff --git a/cpp/include/cuopt/mathematical_optimization/optimization_problem.hpp b/cpp/include/cuopt/mathematical_optimization/optimization_problem.hpp
index 5363cfe812..7d57a50bd8 100644
--- a/cpp/include/cuopt/mathematical_optimization/optimization_problem.hpp
+++ b/cpp/include/cuopt/mathematical_optimization/optimization_problem.hpp
@@ -11,6 +11,7 @@
 #include <cuopt/mathematical_optimization/optimization_problem_interface.hpp>
 #include <cuopt/mathematical_optimization/utilities/internals.hpp>
 
+#include <cuda/stream>
 #include <raft/core/device_span.hpp>
 #include <raft/core/handle.hpp>
 #include <rmm/device_uvector.hpp>
@@ -350,7 +351,7 @@ class optimization_problem_t : public optimization_problem_interface_t<i_t, f_t>
    * @tparam other_f_t  Target floating-point type (e.g. float when this is double)
    */
   template <typename other_f_t>
-  optimization_problem_t<i_t, other_f_t> convert_to_other_prec(rmm::cuda_stream_view stream) const;
+  optimization_problem_t<i_t, other_f_t> convert_to_other_prec(cuda::stream_ref stream) const;
 
   // ============================================================================
   // C API support: Copy to host (polymorphic)
@@ -381,7 +382,7 @@ class optimization_problem_t : public optimization_problem_interface_t<i_t, f_t>
 
  private:
   raft::handle_t const* handle_ptr_{nullptr};
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
 
   problem_category_t problem_category_ = problem_category_t::LP;
   bool maximize_{false};
diff --git a/cpp/include/cuopt/mathematical_optimization/optimization_problem_solution.hpp b/cpp/include/cuopt/mathematical_optimization/optimization_problem_solution.hpp
index b3706473b3..d60982a500 100644
--- a/cpp/include/cuopt/mathematical_optimization/optimization_problem_solution.hpp
+++ b/cpp/include/cuopt/mathematical_optimization/optimization_problem_solution.hpp
@@ -15,7 +15,6 @@
 #include <cuopt/mathematical_optimization/utilities/cython_types.hpp>
 
 #include <raft/core/copy.hpp>
-#include <rmm/cuda_stream_view.hpp>
 
 namespace cuopt {
 namespace CUOPT_EXPORT mathematical_optimization {
diff --git a/cpp/include/cuopt/mathematical_optimization/optimization_problem_solution_interface.hpp b/cpp/include/cuopt/mathematical_optimization/optimization_problem_solution_interface.hpp
index e30474538d..a2548dea16 100644
--- a/cpp/include/cuopt/mathematical_optimization/optimization_problem_solution_interface.hpp
+++ b/cpp/include/cuopt/mathematical_optimization/optimization_problem_solution_interface.hpp
@@ -13,7 +13,6 @@
 #include <cuopt/mathematical_optimization/pdlp/pdlp_warm_start_data.hpp>
 #include <cuopt/mathematical_optimization/pdlp/solver_solution.hpp>  // For pdlp_termination_status_t
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 
 #include <cuopt/mathematical_optimization/utilities/cython_types.hpp>
diff --git a/cpp/include/cuopt/mathematical_optimization/pdlp/pdlp_warm_start_data.hpp b/cpp/include/cuopt/mathematical_optimization/pdlp/pdlp_warm_start_data.hpp
index 52a800c3c2..4d6fd827c2 100644
--- a/cpp/include/cuopt/mathematical_optimization/pdlp/pdlp_warm_start_data.hpp
+++ b/cpp/include/cuopt/mathematical_optimization/pdlp/pdlp_warm_start_data.hpp
@@ -9,6 +9,7 @@
 
 #include <cuopt/export.hpp>
 
+#include <cuda/stream>
 #include <rmm/device_uvector.hpp>
 
 #include <span>
@@ -67,7 +68,7 @@ struct pdlp_warm_start_data_t {
 
   // Copy constructor using the view version for the cython_solver
   pdlp_warm_start_data_t(const pdlp_warm_start_data_view_t<i_t, f_t>& other,
-                         rmm::cuda_stream_view stream_view);
+                         cuda::stream_ref stream_view);
 
   // Copy constructor for when copying the solver_settings object in the PDLP object
   pdlp_warm_start_data_t(const pdlp_warm_start_data_t<i_t, f_t>& other);
diff --git a/cpp/include/cuopt/mathematical_optimization/pdlp/solver_settings.hpp b/cpp/include/cuopt/mathematical_optimization/pdlp/solver_settings.hpp
index 55a3359795..98f2190852 100644
--- a/cpp/include/cuopt/mathematical_optimization/pdlp/solver_settings.hpp
+++ b/cpp/include/cuopt/mathematical_optimization/pdlp/solver_settings.hpp
@@ -8,7 +8,6 @@
 #pragma once
 
 #include <cuopt/mathematical_optimization/constants.h>
-
 #include <cuda/stream>
 #include <cuopt/export.hpp>
 #include <cuopt/mathematical_optimization/cpu_pdlp_warm_start_data.hpp>
@@ -153,7 +152,7 @@ class pdlp_solver_settings_t {
    */
   void set_initial_primal_solution(const f_t* initial_primal_solution,
                                    i_t size,
-                                   rmm::cuda_stream_view stream = cuda::stream_ref{
+                                   cuda::stream_ref stream = cuda::stream_ref{
                                      cudaStream_t{cudaStreamDefault}});
 
   /**
@@ -168,7 +167,7 @@ class pdlp_solver_settings_t {
    */
   void set_initial_dual_solution(const f_t* initial_dual_solution,
                                  i_t size,
-                                 rmm::cuda_stream_view stream = cuda::stream_ref{
+                                 cuda::stream_ref stream = cuda::stream_ref{
                                    cudaStream_t{cudaStreamDefault}});
 
   /** TODO batch mode: tmp
diff --git a/cpp/include/cuopt/mathematical_optimization/pdlp/solver_solution.hpp b/cpp/include/cuopt/mathematical_optimization/pdlp/solver_solution.hpp
index be48ea4baf..0c3e09f274 100644
--- a/cpp/include/cuopt/mathematical_optimization/pdlp/solver_solution.hpp
+++ b/cpp/include/cuopt/mathematical_optimization/pdlp/solver_solution.hpp
@@ -14,7 +14,7 @@
 #include <cuopt/mathematical_optimization/pdlp/solver_settings.hpp>
 #include <cuopt/mathematical_optimization/utilities/internals.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
+#include <cuda/stream>
 #include <rmm/device_uvector.hpp>
 
 #include <raft/core/handle.hpp>
@@ -105,7 +105,7 @@ class optimization_problem_solution_t : public base_solution_t {
    * @param[in] stream_view An rmm view to a stream. All computations will go through this stream
    */
   optimization_problem_solution_t(pdlp_termination_status_t termination_status_,
-                                  rmm::cuda_stream_view stream_view);
+                                  cuda::stream_ref stream_view);
 
   /**
    * @brief Construct an optimization problem solution that serves as PDLP solver output
@@ -115,8 +115,7 @@ class optimization_problem_solution_t : public base_solution_t {
    * 'Optimal', 'PrimalInfeasible', 'DualInfeasible', 'TimeLimit'
    * @param[in] stream_view An rmm view to a stream. All computations will go through this stream
    */
-  optimization_problem_solution_t(cuopt::logic_error error_status_,
-                                  rmm::cuda_stream_view stream_view);
+  optimization_problem_solution_t(cuopt::logic_error error_status_, cuda::stream_ref stream_view);
   /**
    * @brief Construct an optimization problem solution that serves as PDLP solver output
    *
@@ -271,7 +270,7 @@ class optimization_problem_solution_t : public base_solution_t {
    * @param stream_view Non-owning stream view object
    */
   void write_to_file(std::string_view filename,
-                     rmm::cuda_stream_view stream_view,
+                     cuda::stream_ref stream_view,
                      bool generate_variable_values = true);
 
   /**
@@ -281,7 +280,7 @@ class optimization_problem_solution_t : public base_solution_t {
    * @param filename Name of the output file
    * @param stream_view Non-owning stream view object
    */
-  void write_to_sol_file(std::string_view filename, rmm::cuda_stream_view stream_view) const;
+  void write_to_sol_file(std::string_view filename, cuda::stream_ref stream_view) const;
 
   /**
    * @brief Copy solution from another solution object
diff --git a/cpp/include/cuopt/mathematical_optimization/solver_settings.hpp b/cpp/include/cuopt/mathematical_optimization/solver_settings.hpp
index cfe265edcf..0e3a00430e 100644
--- a/cpp/include/cuopt/mathematical_optimization/solver_settings.hpp
+++ b/cpp/include/cuopt/mathematical_optimization/solver_settings.hpp
@@ -11,10 +11,8 @@
 #include <cuopt/mathematical_optimization/pdlp/pdlp_warm_start_data.hpp>
 
 #include <cuda/stream>
-
 #include <raft/core/device_span.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 
 #include <cuopt/mathematical_optimization/constants.h>
@@ -54,11 +52,11 @@ class solver_settings_t {
 
   void set_initial_pdlp_primal_solution(const f_t* initial_primal_solution,
                                         i_t size,
-                                        rmm::cuda_stream_view stream = cuda::stream_ref{
+                                        cuda::stream_ref stream = cuda::stream_ref{
                                           cudaStream_t{cudaStreamDefault}});
   void set_initial_pdlp_dual_solution(const f_t* initial_dual_solution,
                                       i_t size,
-                                      rmm::cuda_stream_view stream = cuda::stream_ref{
+                                      cuda::stream_ref stream = cuda::stream_ref{
                                         cudaStream_t{cudaStreamDefault}});
   void set_pdlp_warm_start_data(const f_t* current_primal_solution,
                                 const f_t* current_dual_solution,
@@ -86,7 +84,7 @@ class solver_settings_t {
   // MIP Settings
   void add_initial_mip_solution(const f_t* initial_solution,
                                 i_t size,
-                                rmm::cuda_stream_view stream = cuda::stream_ref{
+                                cuda::stream_ref stream = cuda::stream_ref{
                                   cudaStream_t{cudaStreamDefault}});
   void set_mip_callback(internals::base_solution_callback_t* callback = nullptr,
                         void* user_data                               = nullptr);
diff --git a/cpp/include/cuopt/mathematical_optimization/utilities/segmented_sum_handler.cuh b/cpp/include/cuopt/mathematical_optimization/utilities/segmented_sum_handler.cuh
index aad0329a7f..6d3d6715dc 100644
--- a/cpp/include/cuopt/mathematical_optimization/utilities/segmented_sum_handler.cuh
+++ b/cpp/include/cuopt/mathematical_optimization/utilities/segmented_sum_handler.cuh
@@ -5,6 +5,7 @@
 
 #pragma once
 
+#include <cuda/stream>
 #include <rmm/device_uvector.hpp>
 
 #include <cub/cub.cuh>
@@ -13,7 +14,7 @@ namespace cuopt {
 
 template <typename i_t, typename f_t>
 struct segmented_sum_handler_t {
-  segmented_sum_handler_t(rmm::cuda_stream_view stream_view) : stream_view_(stream_view) {}
+  segmented_sum_handler_t(cuda::stream_ref stream_view) : stream_view_(stream_view) {}
 
   template <typename InputIteratorT, typename OutputIteratorT>
   void segmented_sum_helper(InputIteratorT input,
@@ -68,7 +69,7 @@ struct segmented_sum_handler_t {
 
   size_t byte_needed_;
   rmm::device_buffer segmented_sum_storage_;
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
 };
 
 }  // namespace cuopt
diff --git a/cpp/include/cuopt/routing/assignment.hpp b/cpp/include/cuopt/routing/assignment.hpp
index b138382d6d..fdd96ad345 100644
--- a/cpp/include/cuopt/routing/assignment.hpp
+++ b/cpp/include/cuopt/routing/assignment.hpp
@@ -7,11 +7,11 @@
 
 #pragma once
 
+#include <cuda/stream>
 #include <cuopt/error.hpp>
 #include <cuopt/export.hpp>
 #include <cuopt/routing/routing_structures.hpp>
 #include <fstream>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 #include <string_view>
 #include <vector>
@@ -54,14 +54,14 @@ class assignment_t {
    * @param status Solution status.
    * @param stream_view Non-owning stream_view object.
    */
-  assignment_t(solution_status_t status, rmm::cuda_stream_view stream_view);
+  assignment_t(solution_status_t status, cuda::stream_ref stream_view);
   /**
    * @brief Constructor.
    *
    * @param error_status Error status.
    * @param stream_view Non-owning stream_view object.
    */
-  assignment_t(cuopt::logic_error error_status, rmm::cuda_stream_view stream_view);
+  assignment_t(cuopt::logic_error error_status, cuda::stream_ref stream_view);
   /**
    * @brief Constructor.
    *
@@ -200,7 +200,7 @@ class assignment_t {
    * @param filename Name of the output file
    * @param stream_view Non-owning stream view object
    */
-  void to_csv(std::string_view filename, rmm::cuda_stream_view stream_view);
+  void to_csv(std::string_view filename, cuda::stream_ref stream_view);
 
   /**
    * @brief Returns the final status as a human readable string
diff --git a/cpp/include/cuopt/routing/distance_engine/waypoint_matrix.hpp b/cpp/include/cuopt/routing/distance_engine/waypoint_matrix.hpp
index e8c04c941d..3b3230cb86 100644
--- a/cpp/include/cuopt/routing/distance_engine/waypoint_matrix.hpp
+++ b/cpp/include/cuopt/routing/distance_engine/waypoint_matrix.hpp
@@ -1,12 +1,13 @@
 /* clang-format off */
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 /* clang-format on */
 
 #pragma once
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
 #include <rmm/device_buffer.hpp>
@@ -168,7 +169,7 @@ class waypoint_matrix_t {
                               f_t const* weights,
                               f_t& out_cost);
   raft::handle_t const* handle_ptr_{nullptr};
-  rmm::cuda_stream_view stream_view_{};
+  cuda::stream_ref stream_view_{};
   i_t const* offsets_;
   i_t n_vertices_;
   i_t const* indices_;
diff --git a/cpp/src/barrier/barrier.cu b/cpp/src/barrier/barrier.cu
index bd55ecfa33..efe06ace9f 100644
--- a/cpp/src/barrier/barrier.cu
+++ b/cpp/src/barrier/barrier.cu
@@ -29,6 +29,7 @@
 
 #include <linear_algebra/vector_math.cuh>
 
+#include <cuda/stream>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -136,7 +137,7 @@ bool should_use_adaptive_regularization(const simplex_solver_settings_t<i_t, f_t
 
 template <typename f_t>
 [[maybe_unused]] static void pairwise_multiply(
-  f_t* a, f_t* b, f_t* out, int size, rmm::cuda_stream_view stream)
+  f_t* a, f_t* b, f_t* out, int size, cuda::stream_ref stream)
 {
   cub::DeviceTransform::Transform(
     cuda::std::make_tuple(a, b), out, size, cuda::std::multiplies<>{}, stream.get());
@@ -145,7 +146,7 @@ template <typename f_t>
 // out[i] = is_direct_free_linear[i] ? 0 : a[i] * b[i]
 template <typename f_t>
 [[maybe_unused]] static void pairwise_multiply_skip_direct_free_linear(
-  f_t* a, f_t* b, int* is_direct_free_linear, f_t* out, int size, rmm::cuda_stream_view stream)
+  f_t* a, f_t* b, int* is_direct_free_linear, f_t* out, int size, cuda::stream_ref stream)
 {
   cub::DeviceTransform::Transform(
     cuda::std::make_tuple(a, b, is_direct_free_linear),
@@ -157,7 +158,7 @@ template <typename f_t>
 
 template <typename f_t>
 [[maybe_unused]] static void axpy(
-  f_t alpha, f_t* x, f_t beta, f_t* y, f_t* out, int size, rmm::cuda_stream_view stream)
+  f_t alpha, f_t* x, f_t beta, f_t* y, f_t* out, int size, cuda::stream_ref stream)
 {
   cub::DeviceTransform::Transform(
     cuda::std::make_tuple(x, y),
@@ -180,7 +181,7 @@ static f2_t<f_t> max_nonnegative_step_length_pair_in_range(
   i_t len,
   const rmm::device_uvector<i_t>& is_direct_free_linear,
   bool apply_direct_free_mask,
-  rmm::cuda_stream_view stream)
+  cuda::stream_ref stream)
 {
   if (len <= 0) { return f2_t<f_t>{f_t(1), f_t(1)}; }
 
@@ -211,7 +212,7 @@ static void recover_linear_orthant_dz(raft::device_span<const f_t> target,
                                       raft::device_span<const f_t> x,
                                       raft::device_span<f_t> dz,
                                       raft::device_span<const i_t> is_direct_free_linear,
-                                      rmm::cuda_stream_view stream)
+                                      cuda::stream_ref stream)
 {
   if (dz.empty()) return;
 
@@ -231,7 +232,7 @@ static void recover_linear_orthant_dz(raft::device_span<const f_t> target,
 template <typename f_t>
 static void negate_complementarity_rhs(raft::device_span<f_t> out,
                                        raft::device_span<const f_t> residual,
-                                       rmm::cuda_stream_view stream)
+                                       cuda::stream_ref stream)
 {
   if (out.empty()) return;
   cub::DeviceTransform::Transform(
@@ -244,7 +245,7 @@ static void fill_linear_cc_rhs(raft::device_span<f_t> out,
                                raft::device_span<const f_t> dz_aff,
                                f_t new_mu,
                                raft::device_span<const i_t> is_direct_free_linear,
-                               rmm::cuda_stream_view stream)
+                               cuda::stream_ref stream)
 {
   if (out.empty()) return;
   cub::DeviceTransform::Transform(
@@ -264,14 +265,14 @@ static void fill_linear_cc_rhs(raft::device_span<f_t> out,
 template <typename i_t, typename f_t>
 class barrier_reduce_helper_t {
  public:
-  explicit barrier_reduce_helper_t(rmm::cuda_stream_view stream_view)
+  explicit barrier_reduce_helper_t(cuda::stream_ref stream_view)
     : d_results_(kCount, stream_view), h_results_(kCount), d_temp_storage_(0, stream_view)
   {
   }
 
   void primal_residual_norm_async(const rmm::device_uvector<f_t>& d_primal_residual,
                                   const rmm::device_uvector<f_t>& d_bound_residual,
-                                  rmm::cuda_stream_view stream_view)
+                                  cuda::stream_ref stream_view)
   {
     norm_inf_async(
       kPrimalResidual, d_primal_residual.data(), d_primal_residual.size(), stream_view);
@@ -279,28 +280,28 @@ class barrier_reduce_helper_t {
   }
 
   void dual_residual_norm_async(const rmm::device_uvector<f_t>& d_dual_residual,
-                                rmm::cuda_stream_view stream_view)
+                                cuda::stream_ref stream_view)
   {
     norm_inf_async(kDualResidual, d_dual_residual.data(), d_dual_residual.size(), stream_view);
   }
 
   void complementarity_residual_norm_async(raft::device_span<const f_t> linear_xz,
                                            const rmm::device_uvector<f_t>& d_wv,
-                                           rmm::cuda_stream_view stream_view)
+                                           cuda::stream_ref stream_view)
   {
     norm_inf_async(kComplXzLinear, linear_xz.data(), linear_xz.size(), stream_view);
     norm_inf_async(kComplWv, d_wv.data(), d_wv.size(), stream_view);
   }
 
   void cone_complementarity_residual_async(raft::device_span<f_t> cone_dot,
-                                           rmm::cuda_stream_view stream_view)
+                                           cuda::stream_ref stream_view)
   {
     max_async(kComplCone, cone_dot.data(), cone_dot.size(), stream_view);
   }
 
   void mu_terms_async(const rmm::device_uvector<f_t>& d_xz,
                       const rmm::device_uvector<f_t>& d_wv,
-                      rmm::cuda_stream_view stream_view)
+                      cuda::stream_ref stream_view)
   {
     sum_async(kMuXzSum, d_xz.data(), d_xz.size(), stream_view);
     sum_async(kMuWvSum, d_wv.data(), d_wv.size(), stream_view);
@@ -309,7 +310,7 @@ class barrier_reduce_helper_t {
   void cTx_async(const rmm::device_uvector<f_t>& d_c,
                  const rmm::device_uvector<f_t>& d_x,
                  cublasHandle_t cublas_handle,
-                 rmm::cuda_stream_view stream_view)
+                 cuda::stream_ref stream_view)
   {
     dot_async(kCTx, d_c, d_x, cublas_handle, stream_view);
   }
@@ -317,7 +318,7 @@ class barrier_reduce_helper_t {
   void bTy_async(const rmm::device_uvector<f_t>& d_b,
                  const rmm::device_uvector<f_t>& d_y,
                  cublasHandle_t cublas_handle,
-                 rmm::cuda_stream_view stream_view)
+                 cuda::stream_ref stream_view)
   {
     dot_async(kBTy, d_b, d_y, cublas_handle, stream_view);
   }
@@ -325,7 +326,7 @@ class barrier_reduce_helper_t {
   void uTv_async(const rmm::device_uvector<f_t>& d_u,
                  const rmm::device_uvector<f_t>& d_v,
                  cublasHandle_t cublas_handle,
-                 rmm::cuda_stream_view stream_view)
+                 cuda::stream_ref stream_view)
   {
     dot_async(kUTv, d_u, d_v, cublas_handle, stream_view);
   }
@@ -333,14 +334,14 @@ class barrier_reduce_helper_t {
   void xTQx_async(const rmm::device_uvector<f_t>& d_Qx,
                   const rmm::device_uvector<f_t>& d_x,
                   cublasHandle_t cublas_handle,
-                  rmm::cuda_stream_view stream_view)
+                  cuda::stream_ref stream_view)
   {
     dot_async(kXTQx, d_Qx, d_x, cublas_handle, stream_view);
   }
 
   // Single batched device-to-host copy + the one stream synchronize needed before any accessor
   // below can be read.
-  void sync(rmm::cuda_stream_view stream_view)
+  void sync(cuda::stream_ref stream_view)
   {
     raft::copy(h_results_.data(), d_results_.data(), static_cast<i_t>(kCount), stream_view);
     stream_view.sync();
@@ -379,7 +380,7 @@ class barrier_reduce_helper_t {
 
   template <typename ReduceOpT>
   void reduce_async(
-    Slot slot, const f_t* in, i_t size, ReduceOpT op, f_t init, rmm::cuda_stream_view stream_view)
+    Slot slot, const f_t* in, i_t size, ReduceOpT op, f_t init, cuda::stream_ref stream_view)
   {
     f_t* out = d_results_.data() + slot;
     if (size == 0) {
@@ -394,17 +395,17 @@ class barrier_reduce_helper_t {
       d_temp_storage_.data(), temp_storage_bytes, in, out, size, op, init, stream_view.get());
   }
 
-  void norm_inf_async(Slot slot, const f_t* in, i_t size, rmm::cuda_stream_view stream_view)
+  void norm_inf_async(Slot slot, const f_t* in, i_t size, cuda::stream_ref stream_view)
   {
     reduce_async(slot, in, size, norm_inf_max{}, f_t(0), stream_view);
   }
 
-  void max_async(Slot slot, const f_t* in, i_t size, rmm::cuda_stream_view stream_view)
+  void max_async(Slot slot, const f_t* in, i_t size, cuda::stream_ref stream_view)
   {
     reduce_async(slot, in, size, thrust::maximum<f_t>{}, f_t(0), stream_view);
   }
 
-  void sum_async(Slot slot, const f_t* in, i_t size, rmm::cuda_stream_view stream_view)
+  void sum_async(Slot slot, const f_t* in, i_t size, cuda::stream_ref stream_view)
   {
     f_t* out                  = d_results_.data() + slot;
     size_t temp_storage_bytes = 0;
@@ -418,7 +419,7 @@ class barrier_reduce_helper_t {
                  const rmm::device_uvector<f_t>& a,
                  const rmm::device_uvector<f_t>& b,
                  cublasHandle_t cublas_handle,
-                 rmm::cuda_stream_view stream_view)
+                 cuda::stream_ref stream_view)
   {
     RAFT_CUBLAS_TRY(raft::linalg::detail::cublasdot(cublas_handle,
                                                     a.size(),
@@ -2294,7 +2295,7 @@ class iteration_data_t {
   bool cone_combined_step_;
   f_t cone_sigma_mu_;
 
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
 
   const simplex_solver_settings_t<i_t, f_t>& settings_;
 };
@@ -3618,7 +3619,7 @@ void fill_linear_complementarity_target(iteration_data_t<i_t, f_t>& data,
                                         raft::device_span<f_t> target,
                                         raft::device_span<const f_t> xz_rhs,
                                         raft::device_span<const f_t> x,
-                                        rmm::cuda_stream_view stream)
+                                        cuda::stream_ref stream)
 {
   if (target.empty()) return;
   cub::DeviceTransform::Transform(
@@ -3637,7 +3638,7 @@ template <typename i_t, typename f_t>
 void fill_affine_cone_complementarity_target(iteration_data_t<i_t, f_t>& data,
                                              i_t cone_var_start,
                                              i_t m_c,
-                                             rmm::cuda_stream_view stream)
+                                             cuda::stream_ref stream)
 {
   if (m_c == 0) return;
   auto& cones = data.cones();
@@ -3656,7 +3657,7 @@ void fill_corrector_cone_complementarity_target(iteration_data_t<i_t, f_t>& data
                                                 i_t cone_var_start,
                                                 i_t m_c,
                                                 f_t sigma_mu,
-                                                rmm::cuda_stream_view stream)
+                                                cuda::stream_ref stream)
 {
   if (m_c == 0) return;
   auto& cones = data.cones();
diff --git a/cpp/src/barrier/barrier.hpp b/cpp/src/barrier/barrier.hpp
index 9865df693a..c1ea9b3403 100644
--- a/cpp/src/barrier/barrier.hpp
+++ b/cpp/src/barrier/barrier.hpp
@@ -17,6 +17,7 @@
 #include <linear_algebra/sparse_matrix.hpp>
 #include <math_optimization/tic_toc.hpp>
 
+#include <cuda/stream>
 #include <rmm/device_uvector.hpp>
 
 #include <utility>
@@ -99,7 +100,7 @@ class barrier_solver_t {
   const simplex::lp_problem_t<i_t, f_t>& lp;
   const simplex::simplex_solver_settings_t<i_t, f_t>& settings;
   const simplex::presolve_info_t<i_t, f_t>& presolve_info;
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
 };
 
 }  // namespace cuopt::mathematical_optimization::barrier
diff --git a/cpp/src/barrier/csr_kkt_build.cuh b/cpp/src/barrier/csr_kkt_build.cuh
index a7667b5012..6ade663c2e 100644
--- a/cpp/src/barrier/csr_kkt_build.cuh
+++ b/cpp/src/barrier/csr_kkt_build.cuh
@@ -9,6 +9,7 @@
 #include <barrier/device_sparse_matrix.cuh>
 #include <barrier/second_order_cone_kernels.cuh>
 
+#include <cuda/stream>
 #include <raft/core/nvtx.hpp>
 #include <utilities/copy_helpers.hpp>
 #include <utilities/cuda_helpers.cuh>
@@ -27,7 +28,7 @@ constexpr int augmented_csr_block_size = 256;
 // Cone -> augmented-KKT-CSR assembly data.
 template <std::integral i_t, std::floating_point f_t>
 struct cone_kkt_data_t {
-  explicit cone_kkt_data_t(rmm::cuda_stream_view stream)
+  explicit cone_kkt_data_t(cuda::stream_ref stream)
     : sparse_ids_by_cone(0, stream),
       dense_ids_by_cone(0, stream),
       dense_cone_entry_rank(0, stream),
@@ -509,7 +510,7 @@ __global__ void fill_augmented_csr_row_kernel(i_t factorization_size,
 template <std::integral i_t, std::floating_point f_t>
 void build_augmented_csr_metadata(const cone_data_t<i_t, f_t>& cones,
                                   cone_kkt_data_t<i_t, f_t>& metadata,
-                                  rmm::cuda_stream_view stream)
+                                  cuda::stream_ref stream)
 {
   raft::common::nvtx::range scope("Barrier: augmented: device CSR metadata");
   const i_t n_cones  = cones.n_cones;
@@ -638,7 +639,7 @@ i_t build_augmented_csr_on_device(i_t n,
                                   cone_kkt_data_t<i_t, f_t>& cone_data,
                                   rmm::device_uvector<i_t>& augmented_diagonal_indices,
                                   device_csr_matrix_t<i_t, f_t>& device_augmented,
-                                  rmm::cuda_stream_view stream)
+                                  cuda::stream_ref stream)
 {
   const i_t factorization_size       = n + m + p;
   const csc_view_t<i_t, f_t> A_view  = A.view();
diff --git a/cpp/src/barrier/device_sparse_matrix.cuh b/cpp/src/barrier/device_sparse_matrix.cuh
index 4012da7a6f..21242ac69d 100644
--- a/cpp/src/barrier/device_sparse_matrix.cuh
+++ b/cpp/src/barrier/device_sparse_matrix.cuh
@@ -11,6 +11,7 @@
 #include <math_optimization/types.hpp>
 
 #include <cub/cub.cuh>
+#include <cuda/stream>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_vector.hpp>
 #include <utilities/copy_helpers.hpp>
@@ -34,13 +35,12 @@ struct sum_reduce_helper_t {
   rmm::device_scalar<f_t> out;
   size_t buffer_size;
 
-  sum_reduce_helper_t(rmm::cuda_stream_view stream_view)
-    : buffer_data(0, stream_view), out(stream_view)
+  sum_reduce_helper_t(cuda::stream_ref stream_view) : buffer_data(0, stream_view), out(stream_view)
   {
   }
 
   template <typename InputIteratorT, typename i_t>
-  f_t sum(InputIteratorT input, i_t size, rmm::cuda_stream_view stream_view)
+  f_t sum(InputIteratorT input, i_t size, cuda::stream_ref stream_view)
   {
     buffer_size = 0;
     cub::DeviceReduce::Sum(nullptr, buffer_size, input, out.data(), size, stream_view.get());
@@ -57,7 +57,7 @@ struct transform_reduce_helper_t {
   rmm::device_scalar<f_t> out;
   size_t buffer_size;
 
-  transform_reduce_helper_t(rmm::cuda_stream_view stream_view)
+  transform_reduce_helper_t(cuda::stream_ref stream_view)
     : buffer_data(0, stream_view), out(stream_view)
   {
   }
@@ -68,7 +68,7 @@ struct transform_reduce_helper_t {
                        TransformOpT transform_op,
                        f_t init,
                        i_t size,
-                       rmm::cuda_stream_view stream_view)
+                       cuda::stream_ref stream_view)
   {
     cub::DeviceReduce::TransformReduce(nullptr,
                                        buffer_size,
@@ -116,7 +116,7 @@ struct transform_reduce_pair_helper_t {
   rmm::device_scalar<f2_t<f_t>> out;
   size_t buffer_size;
 
-  transform_reduce_pair_helper_t(rmm::cuda_stream_view stream_view)
+  transform_reduce_pair_helper_t(cuda::stream_ref stream_view)
     : buffer_data(0, stream_view), out(stream_view)
   {
   }
@@ -128,7 +128,7 @@ struct transform_reduce_pair_helper_t {
                              TransformOpT transform_op,
                              f2_t<f_t> init,
                              i_t size,
-                             rmm::cuda_stream_view stream_view)
+                             cuda::stream_ref stream_view)
   {
     f2_min_t<f_t> reduce_op{};
     cub::DeviceReduce::TransformReduce(nullptr,
@@ -167,12 +167,12 @@ struct csc_view_t {
 template <typename i_t, typename f_t>
 class device_csc_matrix_t {
  public:
-  device_csc_matrix_t(rmm::cuda_stream_view stream)
+  device_csc_matrix_t(cuda::stream_ref stream)
     : col_start(0, stream), i(0, stream), x(0, stream), col_index(0, stream)
   {
   }
 
-  device_csc_matrix_t(i_t rows, i_t cols, i_t nz, rmm::cuda_stream_view stream)
+  device_csc_matrix_t(i_t rows, i_t cols, i_t nz, cuda::stream_ref stream)
     : m(rows),
       n(cols),
       nz_max(nz),
@@ -194,7 +194,7 @@ class device_csc_matrix_t {
   {
   }
 
-  device_csc_matrix_t(const csc_matrix_t<i_t, f_t>& A, rmm::cuda_stream_view stream)
+  device_csc_matrix_t(const csc_matrix_t<i_t, f_t>& A, cuda::stream_ref stream)
     : m(A.m),
       n(A.n),
       nz_max(A.col_start[A.n]),
@@ -208,7 +208,7 @@ class device_csc_matrix_t {
     x         = cuopt::device_copy(A.x, stream);
   }
 
-  void resize_to_nnz(i_t nnz, rmm::cuda_stream_view stream)
+  void resize_to_nnz(i_t nnz, cuda::stream_ref stream)
   {
     col_start.resize(n + 1, stream);
     i.resize(nnz, stream);
@@ -216,7 +216,7 @@ class device_csc_matrix_t {
     nz_max = nnz;
   }
 
-  csc_matrix_t<i_t, f_t> to_host(rmm::cuda_stream_view stream)
+  csc_matrix_t<i_t, f_t> to_host(cuda::stream_ref stream)
   {
     csc_matrix_t<i_t, f_t> A(m, n, nz_max);
     A.col_start = cuopt::host_copy(col_start, stream);
@@ -225,7 +225,7 @@ class device_csc_matrix_t {
     return A;
   }
 
-  void copy(const csc_matrix_t<i_t, f_t>& A, rmm::cuda_stream_view stream)
+  void copy(const csc_matrix_t<i_t, f_t>& A, cuda::stream_ref stream)
   {
     m      = A.m;
     n      = A.n;
@@ -239,7 +239,7 @@ class device_csc_matrix_t {
   }
 
   /** Reset to an empty (all-zero col_start, no nonzeros) matrix of the given shape. */
-  void reset_empty(i_t rows, i_t cols, rmm::cuda_stream_view stream)
+  void reset_empty(i_t rows, i_t cols, cuda::stream_ref stream)
   {
     m      = rows;
     n      = cols;
@@ -250,9 +250,9 @@ class device_csc_matrix_t {
 
   /** Same semantics as csc_matrix_t::to_compressed_row, entirely on
    * device. */
-  void to_compressed_row(device_csr_matrix_t<i_t, f_t>& Arow, rmm::cuda_stream_view stream) const;
+  void to_compressed_row(device_csr_matrix_t<i_t, f_t>& Arow, cuda::stream_ref stream) const;
 
-  void form_col_index(rmm::cuda_stream_view stream)
+  void form_col_index(cuda::stream_ref stream)
   {
     col_index.resize(x.size(), stream);
     RAFT_CUDA_TRY(
@@ -313,12 +313,9 @@ class device_csc_matrix_t {
 template <typename i_t, typename f_t>
 class device_csr_matrix_t {
  public:
-  device_csr_matrix_t(rmm::cuda_stream_view stream)
-    : row_start(0, stream), j(0, stream), x(0, stream)
-  {
-  }
+  device_csr_matrix_t(cuda::stream_ref stream) : row_start(0, stream), j(0, stream), x(0, stream) {}
 
-  device_csr_matrix_t(i_t rows, i_t cols, i_t nz, rmm::cuda_stream_view stream)
+  device_csr_matrix_t(i_t rows, i_t cols, i_t nz, cuda::stream_ref stream)
     : m(rows),
       n(cols),
       nz_max(nz),
@@ -338,7 +335,7 @@ class device_csr_matrix_t {
   {
   }
 
-  device_csr_matrix_t(const csr_matrix_t<i_t, f_t>& A, rmm::cuda_stream_view stream)
+  device_csr_matrix_t(const csr_matrix_t<i_t, f_t>& A, cuda::stream_ref stream)
     : m(A.m),
       n(A.n),
       nz_max(A.row_start[A.m]),
@@ -351,7 +348,7 @@ class device_csr_matrix_t {
     x         = cuopt::device_copy(A.x, stream);
   }
 
-  void resize_to_nnz(i_t nnz, rmm::cuda_stream_view stream)
+  void resize_to_nnz(i_t nnz, cuda::stream_ref stream)
   {
     row_start.resize(m + 1, stream);
     j.resize(nnz, stream);
@@ -359,7 +356,7 @@ class device_csr_matrix_t {
     nz_max = nnz;
   }
 
-  csr_matrix_t<i_t, f_t> to_host(rmm::cuda_stream_view stream)
+  csr_matrix_t<i_t, f_t> to_host(cuda::stream_ref stream)
   {
     csr_matrix_t<i_t, f_t> A(m, n, nz_max);
     A.row_start = cuopt::host_copy(row_start, stream);
@@ -368,7 +365,7 @@ class device_csr_matrix_t {
     return A;
   }
 
-  void copy(csr_matrix_t<i_t, f_t>& A, rmm::cuda_stream_view stream)
+  void copy(csr_matrix_t<i_t, f_t>& A, cuda::stream_ref stream)
   {
     m      = A.m;
     n      = A.n;
@@ -394,7 +391,7 @@ class device_csr_matrix_t {
 
 template <typename i_t, typename f_t>
 void device_csc_matrix_t<i_t, f_t>::to_compressed_row(device_csr_matrix_t<i_t, f_t>& Arow,
-                                                      rmm::cuda_stream_view stream) const
+                                                      cuda::stream_ref stream) const
 {
   static_assert(std::is_signed_v<i_t>);
 
diff --git a/cpp/src/barrier/second_order_cone_kernels.cuh b/cpp/src/barrier/second_order_cone_kernels.cuh
index 16ab9b65e8..7e925d519a 100644
--- a/cpp/src/barrier/second_order_cone_kernels.cuh
+++ b/cpp/src/barrier/second_order_cone_kernels.cuh
@@ -12,6 +12,7 @@
 #include <utilities/copy_helpers.hpp>
 #include <utilities/cuda_helpers.cuh>
 
+#include <cuda/stream>
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
@@ -96,10 +97,7 @@ struct cone_scratch_t {
   // TODO: Consider moving this out to the barrier layer when we wire it in
   rmm::device_uvector<f_t> temp_cone;  // [n_cone_entries]
 
-  cone_scratch_t(i_t n_cones_in,
-                 size_t n_cone_entries_in,
-                 size_t n_large,
-                 rmm::cuda_stream_view stream)
+  cone_scratch_t(i_t n_cones_in, size_t n_cone_entries_in, size_t n_large, cuda::stream_ref stream)
     : n_cones(n_cones_in),
       n_cone_entries(n_cone_entries_in),
       slots(0, stream),
@@ -193,7 +191,7 @@ struct cone_data_t {
   cone_data_t(std::span<const i_t> cone_dimensions_host,
               raft::device_span<f_t> x_in,
               raft::device_span<f_t> z_in,
-              rmm::cuda_stream_view stream,
+              cuda::stream_ref stream,
               i_t soc_threshold_in = 100)
     : n_cones(cone_dimensions_host.size()),
       n_cone_entries(
@@ -446,7 +444,7 @@ __global__ void __launch_bounds__(soc_block_size)
  *   1: ||z_tail||^2 -> z_scale
  */
 template <std::integral i_t, std::floating_point f_t>
-void launch_nt_scaling(cone_data_t<i_t, f_t>& cones, rmm::cuda_stream_view stream)
+void launch_nt_scaling(cone_data_t<i_t, f_t>& cones, cuda::stream_ref stream)
 {
   auto x_scale = cones.scratch.template get_slot<0>();
   auto z_scale = cones.scratch.template get_slot<1>();
@@ -597,7 +595,7 @@ __global__ void update_scaling_sparse_kernel(raft::device_span<const f_t> w,
  * iteration. Call after `launch_nt_scaling` has updated w and eta.
  */
 template <std::integral i_t, std::floating_point f_t>
-void launch_update_scaling_sparse(cone_data_t<i_t, f_t>& cones, rmm::cuda_stream_view stream)
+void launch_update_scaling_sparse(cone_data_t<i_t, f_t>& cones, cuda::stream_ref stream)
 {
   if (!cones.has_sparse_cones()) { return; }
 
@@ -832,7 +830,7 @@ template <std::integral i_t, std::floating_point f_t>
 void apply_w_inv(raft::device_span<const f_t> v,
                  raft::device_span<f_t> out,
                  cone_data_t<i_t, f_t>& cones,
-                 rmm::cuda_stream_view stream)
+                 cuda::stream_ref stream)
 {
   auto w                = cuopt::make_span(cones.w);
   auto eta              = cuopt::make_span(cones.eta);
@@ -868,7 +866,7 @@ template <std::integral i_t, std::floating_point f_t>
 void apply_w(raft::device_span<const f_t> v,
              raft::device_span<f_t> out,
              cone_data_t<i_t, f_t>& cones,
-             rmm::cuda_stream_view stream)
+             cuda::stream_ref stream)
 {
   auto w                = cuopt::make_span(cones.w);
   auto eta              = cuopt::make_span(cones.eta);
@@ -901,7 +899,7 @@ template <std::integral i_t, std::floating_point f_t>
 void apply_hessian(raft::device_span<const f_t> v,
                    raft::device_span<f_t> out,
                    cone_data_t<i_t, f_t>& cones,
-                   rmm::cuda_stream_view stream,
+                   cuda::stream_ref stream,
                    f_t output_scale                  = 1,
                    raft::device_span<const f_t> bias = {},
                    f_t bias_scale                    = 0,
@@ -947,7 +945,7 @@ void recover_cone_dz_from_target(raft::device_span<const f_t> dx,
                                  cone_data_t<i_t, f_t>& cones,
                                  raft::device_span<const f_t> cone_target,
                                  raft::device_span<f_t> dz,
-                                 rmm::cuda_stream_view stream)
+                                 cuda::stream_ref stream)
 {
   apply_hessian<i_t, f_t>(dx, dz, cones, stream, -1, cone_target, 1);
 }
@@ -960,7 +958,7 @@ template <std::integral i_t, std::floating_point f_t>
 void launch_dense_hessian_matvec(raft::device_span<const f_t> x,
                                  cone_data_t<i_t, f_t>& cones,
                                  raft::device_span<f_t> out,
-                                 rmm::cuda_stream_view stream)
+                                 cuda::stream_ref stream)
 {
   auto out_input = raft::device_span<const f_t>(out.data(), out.size());
   apply_hessian<i_t, f_t>(x, out, cones, stream, 1, out_input, 1, true);
@@ -1052,7 +1050,7 @@ void scatter_sparse_hessian_into_augmented(cone_data_t<i_t, f_t>& cones,
                                            const rmm::device_uvector<i_t>& exp_v_row,
                                            const rmm::device_uvector<i_t>& exp_u_row,
                                            const rmm::device_uvector<i_t>& sparse_expansion_D,
-                                           rmm::cuda_stream_view stream,
+                                           cuda::stream_ref stream,
                                            f_t dual_perturb)
 {
   if (!cones.has_sparse_cones()) { return; }
@@ -1169,7 +1167,7 @@ void launch_sparse_augmented_matvec(raft::device_span<const f_t> x,
                                     i_t cone_var_start,
                                     i_t n_primal,
                                     i_t m_constraints,
-                                    rmm::cuda_stream_view stream)
+                                    cuda::stream_ref stream)
 {
   if (!cones.has_sparse_cones()) { return; }
 
@@ -1249,7 +1247,7 @@ void scatter_dense_hessian_into_augmented(const cone_data_t<i_t, f_t>& cones,
                                           const rmm::device_uvector<f_t>& q_values,
                                           const rmm::device_uvector<size_t>& dense_block_offsets,
                                           const rmm::device_uvector<i_t>& dense_cone_ids,
-                                          rmm::cuda_stream_view stream,
+                                          cuda::stream_ref stream,
                                           f_t dual_perturb_value)
 {
   const size_t count = csr_indices.size();
@@ -1458,7 +1456,7 @@ void launch_cone_step_length(segmented_sum_t<i_t>& partitions,
                              raft::device_span<f_t> alpha,
                              raft::device_span<step_tail_sums_t<f_t>> large_sums,
                              f_t alpha_max,
-                             rmm::cuda_stream_view stream)
+                             cuda::stream_ref stream)
 {
   constexpr int warps_per_cta = 8;
   if (!partitions.small_cone_ids.is_empty()) {
@@ -1539,7 +1537,7 @@ f_t compute_cone_step_length(cone_data_t<i_t, f_t>& cones,
                              raft::device_span<const f_t> dx,
                              raft::device_span<const f_t> dz,
                              f_t alpha_max,
-                             rmm::cuda_stream_view stream)
+                             cuda::stream_ref stream)
 {
   auto alpha_primal = cuopt::make_span(cones.scratch.step_alpha_primal);
   auto alpha_dual   = cuopt::make_span(cones.scratch.step_alpha_dual);
@@ -1582,7 +1580,7 @@ void compute_combined_cone_rhs_term(raft::device_span<const f_t> dx_aff,
                                     cone_data_t<i_t, f_t>& cones,
                                     f_t sigma_mu,
                                     raft::device_span<f_t> out,
-                                    rmm::cuda_stream_view stream)
+                                    cuda::stream_ref stream)
 {
   auto cone_offsets     = cuopt::make_span(cones.cone_offsets);
   auto element_cone_ids = cuopt::make_span(cones.element_cone_ids);
diff --git a/cpp/src/barrier/second_order_cone_reduction.cuh b/cpp/src/barrier/second_order_cone_reduction.cuh
index 3e31e5b8be..26627ddef8 100644
--- a/cpp/src/barrier/second_order_cone_reduction.cuh
+++ b/cpp/src/barrier/second_order_cone_reduction.cuh
@@ -14,6 +14,7 @@
 #include <cub/device/device_reduce.cuh>
 #include <cub/warp/warp_reduce.cuh>
 
+#include <cuda/stream>
 #include <rmm/device_buffer.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -83,7 +84,7 @@ struct segmented_sum_t {
 
  private:
   template <typename value_t>
-  void prepare_workspace_for_type(rmm::cuda_stream_view stream)
+  void prepare_workspace_for_type(cuda::stream_ref stream)
   {
     auto input  = thrust::make_constant_iterator(value_t{});
     auto output = thrust::make_discard_iterator();
@@ -106,14 +107,14 @@ struct segmented_sum_t {
 
  public:
   template <typename value_t, typename... rest_t>
-  void prepare_workspace(rmm::cuda_stream_view stream)
+  void prepare_workspace(cuda::stream_ref stream)
   {
     prepare_workspace_for_type<value_t>(stream);
     (prepare_workspace_for_type<rest_t>(stream), ...);
   }
 
   template <typename value_t, typename InputIt, typename OutputIt, int warps_per_cta = 8>
-  void operator()(InputIt input, OutputIt output, value_t init, rmm::cuda_stream_view stream)
+  void operator()(InputIt input, OutputIt output, value_t init, cuda::stream_ref stream)
   {
     if (!small_cone_ids.is_empty()) {
       // Each warp reduces one small cone. `warps_per_cta` only controls how
@@ -153,14 +154,14 @@ struct segmented_sum_t {
   }
 
   template <std::floating_point f_t, typename InputIt>
-  void operator()(InputIt input, raft::device_span<f_t> output, rmm::cuda_stream_view stream)
+  void operator()(InputIt input, raft::device_span<f_t> output, cuda::stream_ref stream)
   {
     operator()(input, output.data(), f_t{0}, stream);
   }
 
   segmented_sum_t(std::span<const i_t> cone_dimensions_host,
                   raft::device_span<const std::size_t> cone_offsets_in,
-                  rmm::cuda_stream_view stream)
+                  cuda::stream_ref stream)
     : cone_offsets(cone_offsets_in),
       small_cone_ids(0, stream),
       medium_cone_ids(0, stream),
diff --git a/cpp/src/barrier/sparse_cholesky.cuh b/cpp/src/barrier/sparse_cholesky.cuh
index dc51cc282d..e73aa0072e 100644
--- a/cpp/src/barrier/sparse_cholesky.cuh
+++ b/cpp/src/barrier/sparse_cholesky.cuh
@@ -16,6 +16,7 @@
 #include <cuda_runtime.h>
 #include <utilities/driver_helpers.cuh>
 
+#include <cuda/stream>
 #include <raft/core/nvtx.hpp>
 
 #include "cudss.h"
@@ -874,7 +875,7 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t<i_t, f_t> {
   bool positive_definite;
   cudaError_t cuda_error;
   cudssStatus_t status;
-  // rmm::cuda_stream_view stream;
+  // cuda::stream_ref stream;
   cudssHandle_t handle;
   cudssDeviceMemHandler_t mem_handler;
   cudssConfig_t solverConfig;
diff --git a/cpp/src/linear_algebra/vector_math.cuh b/cpp/src/linear_algebra/vector_math.cuh
index 85c90c5172..577fd177ac 100644
--- a/cpp/src/linear_algebra/vector_math.cuh
+++ b/cpp/src/linear_algebra/vector_math.cuh
@@ -9,6 +9,7 @@
 
 #include <cub/cub.cuh>
 
+#include <cuda/stream>
 #include <raft/core/copy.hpp>
 #include <raft/core/device_span.hpp>
 #include <raft/core/host_span.hpp>
@@ -37,7 +38,7 @@ struct norm_inf_max {
 };
 
 template <typename i_t, typename f_t, typename InputIteratorT>
-f_t device_custom_vector_norm_inf(InputIteratorT in, i_t size, rmm::cuda_stream_view stream_view)
+f_t device_custom_vector_norm_inf(InputIteratorT in, i_t size, cuda::stream_ref stream_view)
 {
   if (size == 0) { return 0; }
   // FIXME: Tmp storage stored in vector_math class.
@@ -69,13 +70,13 @@ f_t device_custom_vector_norm_inf(InputIteratorT in, i_t size, rmm::cuda_stream_
 }
 
 template <typename i_t, typename f_t>
-f_t device_vector_norm_inf(const rmm::device_uvector<f_t>& in, rmm::cuda_stream_view stream_view)
+f_t device_vector_norm_inf(const rmm::device_uvector<f_t>& in, cuda::stream_ref stream_view)
 {
   return device_custom_vector_norm_inf<i_t, f_t>(in.data(), in.size(), stream_view);
 }
 
 template <typename i_t, typename f_t>
-f_t device_vector_norm_inf(raft::device_span<const f_t> in, rmm::cuda_stream_view stream_view)
+f_t device_vector_norm_inf(raft::device_span<const f_t> in, cuda::stream_ref stream_view)
 {
   return device_custom_vector_norm_inf<i_t, f_t>(in.data(), in.size(), stream_view);
 }
@@ -83,14 +84,14 @@ f_t device_vector_norm_inf(raft::device_span<const f_t> in, rmm::cuda_stream_vie
 // TMP we should just have a CPU and GPU version to do the comparison
 // Should never have to norm inf a CPU vector if we are using the GPU
 template <typename i_t, typename f_t, typename Allocator>
-f_t vector_norm_inf(const std::vector<f_t, Allocator>& x, rmm::cuda_stream_view stream_view)
+f_t vector_norm_inf(const std::vector<f_t, Allocator>& x, cuda::stream_ref stream_view)
 {
   const auto d_x = device_copy(x, stream_view);
   return device_vector_norm_inf<i_t, f_t>(d_x, stream_view);
 }
 
 template <typename i_t, typename f_t>
-f_t vector_norm_inf(raft::host_span<const f_t> x, rmm::cuda_stream_view stream_view)
+f_t vector_norm_inf(raft::host_span<const f_t> x, cuda::stream_ref stream_view)
 {
   rmm::device_uvector<f_t> d_x(x.size(), stream_view);
   raft::copy(d_x.data(), x.data(), x.size(), stream_view);
diff --git a/cpp/src/math_optimization/solver_settings_gpu.cu b/cpp/src/math_optimization/solver_settings_gpu.cu
index a23fbf104a..4ce7261f71 100644
--- a/cpp/src/math_optimization/solver_settings_gpu.cu
+++ b/cpp/src/math_optimization/solver_settings_gpu.cu
@@ -9,7 +9,7 @@
 //
 // Everything else in that class is host-only parameter handling, so the remainder now
 // builds as solver_settings.cpp into the CUDA-free cuopt_client library. Only these
-// members take an rmm::cuda_stream_view or hand back a device_uvector, so they are the
+// members take a cuda::stream_ref or hand back a device_uvector, so they are the
 // only ones that must stay in a CUDA TU inside libcuopt.
 //
 // The `template class` instantiation in solver_settings.cpp cannot emit these members
@@ -17,9 +17,10 @@
 
 #include <cuopt/mathematical_optimization/solver_settings.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 
+#include <cuda/stream>
+
 #include <mip_heuristics/mip_constants.hpp>
 
 namespace cuopt {
@@ -28,7 +29,7 @@ namespace CUOPT_EXPORT mathematical_optimization {
 template <typename i_t, typename f_t>
 void solver_settings_t<i_t, f_t>::set_initial_pdlp_primal_solution(const f_t* solution,
                                                                    i_t size,
-                                                                   rmm::cuda_stream_view stream)
+                                                                   cuda::stream_ref stream)
 {
   pdlp_settings.set_initial_primal_solution(solution, size, stream);
 }
@@ -36,7 +37,7 @@ void solver_settings_t<i_t, f_t>::set_initial_pdlp_primal_solution(const f_t* so
 template <typename i_t, typename f_t>
 void solver_settings_t<i_t, f_t>::set_initial_pdlp_dual_solution(const f_t* solution,
                                                                  i_t size,
-                                                                 rmm::cuda_stream_view stream)
+                                                                 cuda::stream_ref stream)
 {
   pdlp_settings.set_initial_dual_solution(solution, size, stream);
 }
@@ -100,22 +101,22 @@ const rmm::device_uvector<f_t>& solver_settings_t<i_t, f_t>::get_initial_pdlp_du
 template <typename i_t, typename f_t>
 void solver_settings_t<i_t, f_t>::add_initial_mip_solution(const f_t* solution,
                                                            i_t size,
-                                                           rmm::cuda_stream_view stream)
+                                                           cuda::stream_ref stream)
 {
   mip_settings.add_initial_solution(solution, size, stream);
 }
 
 #if MIP_INSTANTIATE_FLOAT
 template CUOPT_EXPORT void solver_settings_t<int, float>::set_initial_pdlp_primal_solution(
-  const float*, int, rmm::cuda_stream_view);
+  const float*, int, cuda::stream_ref);
 template CUOPT_EXPORT void solver_settings_t<int, float>::set_initial_pdlp_dual_solution(
-  const float*, int, rmm::cuda_stream_view);
+  const float*, int, cuda::stream_ref);
 template CUOPT_EXPORT const rmm::device_uvector<float>&
 solver_settings_t<int, float>::get_initial_pdlp_primal_solution() const;
 template CUOPT_EXPORT const rmm::device_uvector<float>&
 solver_settings_t<int, float>::get_initial_pdlp_dual_solution() const;
 template CUOPT_EXPORT void solver_settings_t<int, float>::add_initial_mip_solution(
-  const float*, int, rmm::cuda_stream_view);
+  const float*, int, cuda::stream_ref);
 // The 19-argument host overload. It was moved into this TU with the rest of the block, but
 // `template class` in solver_settings.cpp cannot emit it (definition not visible there), so
 // without this line the symbol disappears -- and it is the one the Cython layer binds to,
@@ -143,15 +144,15 @@ template CUOPT_EXPORT void solver_settings_t<int, float>::set_pdlp_warm_start_da
 
 #if MIP_INSTANTIATE_DOUBLE
 template CUOPT_EXPORT void solver_settings_t<int, double>::set_initial_pdlp_primal_solution(
-  const double*, int, rmm::cuda_stream_view);
+  const double*, int, cuda::stream_ref);
 template CUOPT_EXPORT void solver_settings_t<int, double>::set_initial_pdlp_dual_solution(
-  const double*, int, rmm::cuda_stream_view);
+  const double*, int, cuda::stream_ref);
 template CUOPT_EXPORT const rmm::device_uvector<double>&
 solver_settings_t<int, double>::get_initial_pdlp_primal_solution() const;
 template CUOPT_EXPORT const rmm::device_uvector<double>&
 solver_settings_t<int, double>::get_initial_pdlp_dual_solution() const;
 template CUOPT_EXPORT void solver_settings_t<int, double>::add_initial_mip_solution(
-  const double*, int, rmm::cuda_stream_view);
+  const double*, int, cuda::stream_ref);
 // The 19-argument host overload. It was moved into this TU with the rest of the block, but
 // `template class` in solver_settings.cpp cannot emit it (definition not visible there), so
 // without this line the symbol disappears -- and it is the one the Cython layer binds to,
diff --git a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu
index 0fa6b3c3d3..92fbfb930c 100644
--- a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu
+++ b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu
@@ -17,6 +17,7 @@
 #include <utilities/device_scalar_init.hpp>
 #include <utilities/timer.hpp>
 
+#include <cuda/stream>
 #include <raft/linalg/eltwise.cuh>
 #include <raft/linalg/reduce.cuh>
 #include <raft/random/rng.cuh>
@@ -122,7 +123,7 @@ fj_t<i_t, f_t>::~fj_t()
 }
 
 template <typename i_t, typename f_t>
-void fj_t<i_t, f_t>::reset_weights(const rmm::cuda_stream_view& climber_stream, f_t weight)
+void fj_t<i_t, f_t>::reset_weights(const cuda::stream_ref& climber_stream, f_t weight)
 {
   // unless reset explicitly, the values are kept across runs and across climbers
   max_cstr_weight.set_value_async(weight, climber_stream);
@@ -280,7 +281,7 @@ void fj_t<i_t, f_t>::copy_weights(const weight_t<i_t, f_t>& weights,
 }
 
 template <typename i_t, typename f_t>
-void fj_t<i_t, f_t>::climber_data_t::clear_sets(const rmm::cuda_stream_view& stream)
+void fj_t<i_t, f_t>::climber_data_t::clear_sets(const cuda::stream_ref& stream)
 {
   violated_constraints.clear(stream);
   candidate_variables.clear(stream);
@@ -289,7 +290,7 @@ void fj_t<i_t, f_t>::climber_data_t::clear_sets(const rmm::cuda_stream_view& str
 }
 
 template <typename i_t, typename f_t>
-void fj_t<i_t, f_t>::device_init(const rmm::cuda_stream_view& stream)
+void fj_t<i_t, f_t>::device_init(const cuda::stream_ref& stream)
 {
   thrust::for_each(rmm::exec_policy(stream),
                    thrust::counting_iterator<i_t>(0),
@@ -317,7 +318,7 @@ void fj_t<i_t, f_t>::climber_init(i_t climber_idx)
 }
 
 template <typename i_t, typename f_t>
-void fj_t<i_t, f_t>::climber_init(i_t climber_idx, const rmm::cuda_stream_view& climber_stream)
+void fj_t<i_t, f_t>::climber_init(i_t climber_idx, const cuda::stream_ref& climber_stream)
 {
   raft::common::nvtx::range scope("climber_init");
 
@@ -603,8 +604,7 @@ void fj_t<i_t, f_t>::run_step_device(i_t climber_idx, bool use_graph)
 
 // TODO: switch to conditional graph nodes once we switch to CTK >= 12.4
 template <typename i_t, typename f_t>
-void fj_t<i_t, f_t>::load_balancing_score_update(const rmm::cuda_stream_view& stream,
-                                                 i_t climber_idx)
+void fj_t<i_t, f_t>::load_balancing_score_update(const cuda::stream_ref& stream, i_t climber_idx)
 {
   auto [grid_load_balancing_prepare, blocks_load_balancing_prepare] =
     load_balancing_prepare_launch_dims;
@@ -660,7 +660,7 @@ void fj_t<i_t, f_t>::load_balancing_score_update(const rmm::cuda_stream_view& st
 }
 
 template <typename i_t, typename f_t>
-void fj_t<i_t, f_t>::run_step_device(const rmm::cuda_stream_view& climber_stream,
+void fj_t<i_t, f_t>::run_step_device(const cuda::stream_ref& climber_stream,
                                      i_t climber_idx,
                                      bool use_graph)
 {
@@ -811,7 +811,7 @@ void fj_t<i_t, f_t>::round_remaining_fractionals(solution_t<i_t, f_t>& solution,
 }
 
 template <typename i_t, typename f_t>
-void fj_t<i_t, f_t>::refresh_lhs_and_violation(const rmm::cuda_stream_view& stream, i_t climber_idx)
+void fj_t<i_t, f_t>::refresh_lhs_and_violation(const cuda::stream_ref& stream, i_t climber_idx)
 {
   auto& data = *climbers[climber_idx];
   auto v     = data.view();
diff --git a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cuh b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cuh
index a0f3103233..6f7550c0b6 100644
--- a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cuh
+++ b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cuh
@@ -7,6 +7,7 @@
 
 #pragma once
 
+#include <cuda/stream>
 #include "utils.cuh"
 
 #include <cuopt/mathematical_optimization/mip/solver_settings.hpp>
@@ -232,23 +233,21 @@ class fj_t {
     bool randomize_params  = false);
   i_t alloc_max_climbers(i_t desired_climbers);
   void resize_vectors(const raft::handle_t* handle_ptr);
-  void device_init(const rmm::cuda_stream_view& stream);
+  void device_init(const cuda::stream_ref& stream);
   void climber_init(i_t climber_idx);
-  void climber_init(i_t climber_idx, const rmm::cuda_stream_view& stream);
+  void climber_init(i_t climber_idx, const cuda::stream_ref& stream);
   void set_fj_settings(fj_settings_t settings_);
-  void reset_weights(const rmm::cuda_stream_view& stream, f_t weight = 10.);
+  void reset_weights(const cuda::stream_ref& stream, f_t weight = 10.);
   void randomize_weights(const raft::handle_t* handle_ptr);
   void copy_weights(const weight_t<i_t, f_t>& weights,
                     const raft::handle_t* handle_ptr,
                     std::optional<i_t> new_size = std::nullopt);
   i_t host_loop(solution_t<i_t, f_t>& solution, i_t climber_idx = 0);
   void run_step_device(i_t climber_idx = 0, bool use_graph = true);
-  void run_step_device(const rmm::cuda_stream_view& stream,
-                       i_t climber_idx = 0,
-                       bool use_graph  = true);
-  void refresh_lhs_and_violation(const rmm::cuda_stream_view& stream, i_t climber_idx = 0);
+  void run_step_device(const cuda::stream_ref& stream, i_t climber_idx = 0, bool use_graph = true);
+  void refresh_lhs_and_violation(const cuda::stream_ref& stream, i_t climber_idx = 0);
   // load balancing
-  void load_balancing_score_update(const rmm::cuda_stream_view& stream, i_t climber_idx = 0);
+  void load_balancing_score_update(const cuda::stream_ref& stream, i_t climber_idx = 0);
   // executed after a roudning FJ run if any fractionals remain to eliminate them
   void round_remaining_fractionals(solution_t<i_t, f_t>& solution, i_t climber_idx = 0);
 
@@ -642,7 +641,7 @@ class fj_t {
     };
 
     view_t view();
-    void clear_sets(const rmm::cuda_stream_view& stream);
+    void clear_sets(const cuda::stream_ref& stream);
   };
   void populate_climber_views();
 
diff --git a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cu b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cu
index 0469574197..2f80107823 100644
--- a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cu
+++ b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cu
@@ -12,6 +12,7 @@
 #include <mip_heuristics/logger.cuh>
 #include <utilities/device_utils.cuh>
 
+#include <cuda/stream>
 #include <raft/random/rng.cuh>
 
 #include <thrust/iterator/transform_iterator.h>
@@ -1439,7 +1440,7 @@ template <typename i_t, typename f_t>
 void launch_load_balancing_prepare_iteration(dim3 grid,
                                              dim3 blocks,
                                              void** kernel_args,
-                                             rmm::cuda_stream_view stream)
+                                             cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchCooperativeKernel(
     (void*)load_balancing_prepare_iteration<i_t, f_t>, grid, blocks, kernel_args, 0, stream.get()));
@@ -1457,7 +1458,7 @@ template <typename i_t, typename f_t>
 void launch_update_assignment_kernel(dim3 grid,
                                      dim3 blocks,
                                      void** kernel_args,
-                                     rmm::cuda_stream_view stream)
+                                     cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchKernel(
     (void*)update_assignment_kernel<i_t, f_t>, grid, blocks, kernel_args, 0, stream.get()));
@@ -1531,7 +1532,7 @@ template <typename i_t, typename f_t, MTMMoveType move_type, bool is_binary_pb>
 void launch_compute_mtm_moves_kernel(dim3 grid,
                                      dim3 blocks,
                                      void** kernel_args,
-                                     rmm::cuda_stream_view stream)
+                                     cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(
     cudaLaunchCooperativeKernel((void*)compute_mtm_moves_kernel<i_t, f_t, move_type, is_binary_pb>,
@@ -1546,7 +1547,7 @@ template <typename i_t, typename f_t>
 void launch_load_balancing_sanity_checks(dim3 grid,
                                          dim3 blocks,
                                          void** kernel_args,
-                                         rmm::cuda_stream_view stream)
+                                         cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchCooperativeKernel(
     (void*)load_balancing_sanity_checks<i_t, f_t>, grid, blocks, kernel_args, 0, stream.get()));
@@ -1556,7 +1557,7 @@ template <typename i_t, typename f_t>
 void launch_handle_local_minimum_kernel(dim3 grid,
                                         dim3 blocks,
                                         void** kernel_args,
-                                        rmm::cuda_stream_view stream)
+                                        cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchCooperativeKernel(
     (void*)handle_local_minimum_kernel<i_t, f_t>, grid, blocks, kernel_args, 0, stream.get()));
@@ -1574,7 +1575,7 @@ template <typename i_t, typename f_t>
 void launch_update_changed_constraints_kernel(dim3 grid,
                                               dim3 blocks,
                                               void** kernel_args,
-                                              rmm::cuda_stream_view stream)
+                                              cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchKernel((void*)update_changed_constraints_kernel<i_t, f_t>,
                                  grid,
@@ -1588,7 +1589,7 @@ template <typename i_t, typename f_t>
 void launch_update_lift_moves_kernel(dim3 grid,
                                      dim3 blocks,
                                      void** kernel_args,
-                                     rmm::cuda_stream_view stream)
+                                     cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchKernel(
     (void*)update_lift_moves_kernel<i_t, f_t>, grid, blocks, kernel_args, 0, stream.get()));
@@ -1598,7 +1599,7 @@ template <typename i_t, typename f_t>
 void launch_update_breakthrough_moves_kernel(dim3 grid,
                                              dim3 blocks,
                                              void** kernel_args,
-                                             rmm::cuda_stream_view stream)
+                                             cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchKernel(
     (void*)update_breakthrough_moves_kernel<i_t, f_t>, grid, blocks, kernel_args, 0, stream.get()));
@@ -1608,7 +1609,7 @@ template <typename i_t, typename f_t>
 void launch_select_variable_kernel(dim3 grid,
                                    dim3 blocks,
                                    void** kernel_args,
-                                   rmm::cuda_stream_view stream)
+                                   cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchKernel(
     (void*)select_variable_kernel<i_t, f_t>, grid, blocks, kernel_args, 0, stream.get()));
@@ -1618,7 +1619,7 @@ template <typename i_t, typename f_t>
 void launch_init_lhs_and_violation(dim3 grid,
                                    dim3 blocks,
                                    void** kernel_args,
-                                   rmm::cuda_stream_view stream)
+                                   cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchKernel(
     (void*)init_lhs_and_violation<i_t, f_t>, grid, blocks, kernel_args, 0, stream.get()));
@@ -1628,7 +1629,7 @@ template <typename i_t, typename f_t>
 void launch_update_best_solution_kernel(dim3 grid,
                                         dim3 blocks,
                                         void** kernel_args,
-                                        rmm::cuda_stream_view stream)
+                                        cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchKernel(
     (void*)update_best_solution_kernel<i_t, f_t>, grid, blocks, kernel_args, 0, stream.get()));
@@ -1638,7 +1639,7 @@ template <typename i_t, typename f_t>
 void launch_load_balancing_compute_workid_mappings(dim3 grid,
                                                    dim3 blocks,
                                                    void** kernel_args,
-                                                   rmm::cuda_stream_view stream)
+                                                   cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchKernel((void*)load_balancing_compute_workid_mappings<i_t, f_t>,
                                  grid,
@@ -1652,7 +1653,7 @@ template <typename i_t, typename f_t>
 void launch_load_balancing_init_cstr_bounds_csr(dim3 grid,
                                                 dim3 blocks,
                                                 void** kernel_args,
-                                                rmm::cuda_stream_view stream)
+                                                cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchKernel((void*)load_balancing_init_cstr_bounds_csr<i_t, f_t>,
                                  grid,
@@ -1666,7 +1667,7 @@ template <typename i_t, typename f_t>
 void launch_load_balancing_compute_scores_binary(dim3 grid,
                                                  dim3 blocks,
                                                  void** kernel_args,
-                                                 rmm::cuda_stream_view stream)
+                                                 cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchKernel((void*)load_balancing_compute_scores_binary<i_t, f_t>,
                                  grid,
@@ -1680,7 +1681,7 @@ template <typename i_t, typename f_t>
 void launch_load_balancing_mtm_compute_candidates(dim3 grid,
                                                   dim3 blocks,
                                                   void** kernel_args,
-                                                  rmm::cuda_stream_view stream)
+                                                  cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchKernel((void*)load_balancing_mtm_compute_candidates<i_t, f_t>,
                                  grid,
@@ -1694,7 +1695,7 @@ template <typename i_t, typename f_t>
 void launch_load_balancing_mtm_compute_scores(dim3 grid,
                                               dim3 blocks,
                                               void** kernel_args,
-                                              rmm::cuda_stream_view stream)
+                                              cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchKernel((void*)load_balancing_mtm_compute_scores<i_t, f_t>,
                                  grid,
@@ -1758,11 +1759,11 @@ void launch_load_balancing_mtm_compute_scores(dim3 grid,
   template __global__ void select_variable_kernel<int, F_TYPE>(                                    \
     typename fj_t<int, F_TYPE>::climber_data_t::view_t fj);                                        \
   template void launch_load_balancing_prepare_iteration<int, F_TYPE>(                              \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template std::pair<dim3, dim3> get_launch_dims_update_assignment_kernel<int, F_TYPE>(            \
     int TPB, const raft::handle_t* handle_ptr);                                                    \
   template void launch_update_assignment_kernel<int, F_TYPE>(                                      \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template std::pair<dim3, dim3>                                                                   \
   get_launch_dims_compute_mtm_moves_kernel<int, F_TYPE, MTMMoveType::FJ_MTM_VIOLATED, false>(      \
     int TPB, const raft::handle_t* handle_ptr);                                                    \
@@ -1787,37 +1788,37 @@ void launch_load_balancing_mtm_compute_scores(dim3 grid,
   template std::pair<dim3, dim3> get_launch_dims_load_balancing_prepare_iteration<int, F_TYPE>(    \
     int TPB, const raft::handle_t* handle_ptr);                                                    \
   template void launch_compute_mtm_moves_kernel<int, F_TYPE, MTMMoveType::FJ_MTM_VIOLATED, false>( \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template void launch_compute_mtm_moves_kernel<int, F_TYPE, MTMMoveType::FJ_MTM_VIOLATED, true>(  \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template void launch_load_balancing_sanity_checks<int, F_TYPE>(                                  \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template void launch_handle_local_minimum_kernel<int, F_TYPE>(                                   \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template std::pair<dim3, dim3> get_launch_dims_update_changed_constraints_kernel<int, F_TYPE>(   \
     int TPB, const raft::handle_t* handle_ptr);                                                    \
   template void launch_update_changed_constraints_kernel<int, F_TYPE>(                             \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template void launch_update_lift_moves_kernel<int, F_TYPE>(                                      \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template void launch_update_breakthrough_moves_kernel<int, F_TYPE>(                              \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template void launch_select_variable_kernel<int, F_TYPE>(                                        \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template void launch_init_lhs_and_violation<int, F_TYPE>(                                        \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template void launch_update_best_solution_kernel<int, F_TYPE>(                                   \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template void launch_load_balancing_compute_workid_mappings<int, F_TYPE>(                        \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template void launch_load_balancing_init_cstr_bounds_csr<int, F_TYPE>(                           \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template void launch_load_balancing_compute_scores_binary<int, F_TYPE>(                          \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template void launch_load_balancing_mtm_compute_candidates<int, F_TYPE>(                         \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);                     \
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);                          \
   template void launch_load_balancing_mtm_compute_scores<int, F_TYPE>(                             \
-    dim3 grid, dim3 blocks, void** kernel_args, rmm::cuda_stream_view stream);
+    dim3 grid, dim3 blocks, void** kernel_args, cuda::stream_ref stream);
 
 #if MIP_INSTANTIATE_FLOAT
 CUOPT_INSTANTIATE(float)
diff --git a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cuh b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cuh
index 3bc4b2ebaa..faaa4f6227 100644
--- a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cuh
+++ b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump_kernels.cuh
@@ -12,6 +12,7 @@
 #include <mip_heuristics/logger.cuh>
 #include <utilities/device_utils.cuh>
 
+#include <cuda/stream>
 #include <raft/random/rng.cuh>
 
 #include <cub/cub.cuh>
@@ -102,7 +103,7 @@ template <typename i_t, typename f_t>
 void launch_load_balancing_prepare_iteration(dim3 grid,
                                              dim3 blocks,
                                              void** kernel_args,
-                                             rmm::cuda_stream_view stream);
+                                             cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 std::pair<dim3, dim3> get_launch_dims_update_assignment_kernel(int TPB,
@@ -112,7 +113,7 @@ template <typename i_t, typename f_t>
 void launch_update_assignment_kernel(dim3 grid,
                                      dim3 blocks,
                                      void** kernel_args,
-                                     rmm::cuda_stream_view stream);
+                                     cuda::stream_ref stream);
 
 template <typename i_t, typename f_t, MTMMoveType move_type, bool is_binary_pb>
 std::pair<dim3, dim3> get_launch_dims_compute_mtm_moves_kernel(int TPB,
@@ -150,19 +151,19 @@ template <typename i_t, typename f_t, MTMMoveType move_type, bool is_binary_pb>
 void launch_compute_mtm_moves_kernel(dim3 grid,
                                      dim3 blocks,
                                      void** kernel_args,
-                                     rmm::cuda_stream_view stream);
+                                     cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 void launch_load_balancing_sanity_checks(dim3 grid,
                                          dim3 blocks,
                                          void** kernel_args,
-                                         rmm::cuda_stream_view stream);
+                                         cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 void launch_handle_local_minimum_kernel(dim3 grid,
                                         dim3 blocks,
                                         void** kernel_args,
-                                        rmm::cuda_stream_view stream);
+                                        cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 std::pair<dim3, dim3> get_launch_dims_update_changed_constraints_kernel(
@@ -172,66 +173,66 @@ template <typename i_t, typename f_t>
 void launch_update_changed_constraints_kernel(dim3 grid,
                                               dim3 blocks,
                                               void** kernel_args,
-                                              rmm::cuda_stream_view stream);
+                                              cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 void launch_update_lift_moves_kernel(dim3 grid,
                                      dim3 blocks,
                                      void** kernel_args,
-                                     rmm::cuda_stream_view stream);
+                                     cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 void launch_update_breakthrough_moves_kernel(dim3 grid,
                                              dim3 blocks,
                                              void** kernel_args,
-                                             rmm::cuda_stream_view stream);
+                                             cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 void launch_select_variable_kernel(dim3 grid,
                                    dim3 blocks,
                                    void** kernel_args,
-                                   rmm::cuda_stream_view stream);
+                                   cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 void launch_init_lhs_and_violation(dim3 grid,
                                    dim3 blocks,
                                    void** kernel_args,
-                                   rmm::cuda_stream_view stream);
+                                   cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 void launch_update_best_solution_kernel(dim3 grid,
                                         dim3 blocks,
                                         void** kernel_args,
-                                        rmm::cuda_stream_view stream);
+                                        cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 void launch_load_balancing_compute_workid_mappings(dim3 grid,
                                                    dim3 blocks,
                                                    void** kernel_args,
-                                                   rmm::cuda_stream_view stream);
+                                                   cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 void launch_load_balancing_init_cstr_bounds_csr(dim3 grid,
                                                 dim3 blocks,
                                                 void** kernel_args,
-                                                rmm::cuda_stream_view stream);
+                                                cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 void launch_load_balancing_compute_scores_binary(dim3 grid,
                                                  dim3 blocks,
                                                  void** kernel_args,
-                                                 rmm::cuda_stream_view stream);
+                                                 cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 void launch_load_balancing_mtm_compute_candidates(dim3 grid,
                                                   dim3 blocks,
                                                   void** kernel_args,
-                                                  rmm::cuda_stream_view stream);
+                                                  cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 void launch_load_balancing_mtm_compute_scores(dim3 grid,
                                               dim3 blocks,
                                               void** kernel_args,
-                                              rmm::cuda_stream_view stream);
+                                              cuda::stream_ref stream);
 
 }  // namespace cuopt::mathematical_optimization::mip
diff --git a/cpp/src/mip_heuristics/feasibility_jump/utils.cuh b/cpp/src/mip_heuristics/feasibility_jump/utils.cuh
index 7eee62e8a3..1fb3c1854c 100644
--- a/cpp/src/mip_heuristics/feasibility_jump/utils.cuh
+++ b/cpp/src/mip_heuristics/feasibility_jump/utils.cuh
@@ -11,6 +11,7 @@
 
 #include <thrust/pair.h>
 #include <cuda/atomic>
+#include <cuda/stream>
 #include <raft/core/device_span.hpp>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
@@ -36,13 +37,13 @@ template <typename word_t = uint32_t>
 struct bitmap_t {
   static constexpr int bits_per_word = sizeof(word_t) * CHAR_BIT;
 
-  bitmap_t(size_t size, const rmm::cuda_stream_view& stream)
+  bitmap_t(size_t size, const cuda::stream_ref& stream)
     : validity_bitmap(size > 0 ? (size - 1) / bits_per_word + 1 : 0, stream)
   {
     clear(stream);
   }
 
-  void clear(const rmm::cuda_stream_view& stream)
+  void clear(const cuda::stream_ref& stream)
   {
     cudaMemsetAsync(
       validity_bitmap.data(), 0, sizeof(word_t) * validity_bitmap.size(), stream.get());
@@ -52,7 +53,7 @@ struct bitmap_t {
     thrust::uninitialized_fill(
       handle_ptr->get_thrust_policy(), validity_bitmap.begin(), validity_bitmap.end(), 0);
   }
-  void resize(size_t size, const rmm::cuda_stream_view& stream)
+  void resize(size_t size, const cuda::stream_ref& stream)
   {
     validity_bitmap.resize(size > 0 ? (size - 1) / bits_per_word + 1 : 0, stream);
   }
@@ -100,7 +101,7 @@ struct bitmap_t {
 
 template <typename i_t, typename f_t>
 struct contiguous_set_t {
-  contiguous_set_t(i_t max_size, const rmm::cuda_stream_view& stream)
+  contiguous_set_t(i_t max_size, const cuda::stream_ref& stream)
     : set_size(zero_v<i_t>, stream),
       lock(zero_v<i_t>, stream),
       contents(max_size, stream),
@@ -110,7 +111,7 @@ struct contiguous_set_t {
     clear(stream);
   }
 
-  void clear(const rmm::cuda_stream_view& stream)
+  void clear(const cuda::stream_ref& stream)
   {
     set_size.set_value_to_zero_async(stream);
     // can't use thrust::fill, needs a memset node in order to be recorded in CUDA graphs
@@ -127,7 +128,7 @@ struct contiguous_set_t {
     set_size.set_value_to_zero_async(handle_ptr->get_stream());
   }
 
-  void resize(size_t size, const rmm::cuda_stream_view& stream)
+  void resize(size_t size, const cuda::stream_ref& stream)
   {
     contents.resize(size, stream);
     index_map.resize(size, stream);
diff --git a/cpp/src/mip_heuristics/mip_scaling_strategy.cu b/cpp/src/mip_heuristics/mip_scaling_strategy.cu
index 2ab03ac535..41525a9c65 100644
--- a/cpp/src/mip_heuristics/mip_scaling_strategy.cu
+++ b/cpp/src/mip_heuristics/mip_scaling_strategy.cu
@@ -10,6 +10,7 @@
 #include <pdlp/utils.cuh>
 #include <utilities/logger.hpp>
 
+#include <cuda/stream>
 #include <raft/util/cudart_utils.hpp>
 
 #include <cub/cub.cuh>
@@ -149,7 +150,7 @@ void compute_row_inf_norm(
   rmm::device_uvector<std::uint8_t>& temp_storage,
   size_t temp_storage_bytes,
   rmm::device_uvector<f_t>& row_inf_norm,
-  rmm::cuda_stream_view stream_view)
+  cuda::stream_ref stream_view)
 {
   const auto& matrix_values  = op_problem.get_constraint_matrix_values();
   const auto& matrix_offsets = op_problem.get_constraint_matrix_offsets();
@@ -174,7 +175,7 @@ void compute_row_integer_gcd(
   rmm::device_uvector<std::uint8_t>& temp_storage,
   size_t temp_storage_bytes,
   rmm::device_uvector<std::int64_t>& row_integer_gcd,
-  rmm::cuda_stream_view stream_view)
+  cuda::stream_ref stream_view)
 {
   const auto& matrix_values  = op_problem.get_constraint_matrix_values();
   const auto& matrix_indices = op_problem.get_constraint_matrix_indices();
@@ -357,7 +358,7 @@ rmm::device_uvector<std::int64_t> capture_pre_scaling_integer_gcd(
   const cuopt::mathematical_optimization::optimization_problem_t<i_t, f_t>& op_problem,
   rmm::device_uvector<std::uint8_t>& temp_storage,
   size_t temp_storage_bytes,
-  rmm::cuda_stream_view stream_view)
+  cuda::stream_ref stream_view)
 {
   const i_t n_rows = op_problem.get_n_constraints();
   rmm::device_uvector<std::int64_t> gcd(static_cast<size_t>(n_rows), stream_view);
@@ -371,7 +372,7 @@ void assert_integer_coefficient_integrality(
   rmm::device_uvector<std::uint8_t>& temp_storage,
   size_t temp_storage_bytes,
   const rmm::device_uvector<std::int64_t>& pre_scaling_gcd,
-  rmm::cuda_stream_view stream_view)
+  cuda::stream_ref stream_view)
 {
   const auto* handle_ptr = op_problem.get_handle_ptr();
   const i_t n_rows       = op_problem.get_n_constraints();
@@ -415,7 +416,7 @@ size_t dry_run_cub(
   rmm::device_uvector<f_t>& row_min_nonzero,
   rmm::device_uvector<i_t>& row_nonzero_count,
   rmm::device_uvector<std::int64_t>& row_integer_gcd,
-  rmm::cuda_stream_view stream_view)
+  cuda::stream_ref stream_view)
 {
   const auto& matrix_values     = op_problem.get_constraint_matrix_values();
   const auto& matrix_indices    = op_problem.get_constraint_matrix_indices();
diff --git a/cpp/src/mip_heuristics/mip_scaling_strategy.cuh b/cpp/src/mip_heuristics/mip_scaling_strategy.cuh
index 3355eee5ca..3f2796c128 100644
--- a/cpp/src/mip_heuristics/mip_scaling_strategy.cuh
+++ b/cpp/src/mip_heuristics/mip_scaling_strategy.cuh
@@ -9,10 +9,9 @@
 
 #include <cuopt/mathematical_optimization/optimization_problem.hpp>
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
-
 namespace cuopt::mathematical_optimization::mip {
 
 template <typename i_t, typename f_t>
@@ -26,7 +25,7 @@ class mip_scaling_strategy_t {
 
  private:
   raft::handle_t const* handle_ptr_{nullptr};
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
   optimization_problem_type_t& op_problem_scaled_;
 };
 
diff --git a/cpp/src/mip_heuristics/presolve/conditional_bound_strengthening.cu b/cpp/src/mip_heuristics/presolve/conditional_bound_strengthening.cu
index 0223c77422..14258e6b16 100644
--- a/cpp/src/mip_heuristics/presolve/conditional_bound_strengthening.cu
+++ b/cpp/src/mip_heuristics/presolve/conditional_bound_strengthening.cu
@@ -12,6 +12,7 @@
 #include <utilities/vector_helpers.cuh>
 
 #include <raft/sparse/detail/cusparse_wrappers.h>
+#include <cuda/stream>
 #include <raft/core/cusparse_macros.hpp>
 #include <raft/sparse/linalg/transpose.cuh>
 #include "cusparse.h"
@@ -642,7 +643,7 @@ struct len_from_offset {
 // Ideally this should be precomputed and stored in the problem, but that also means we need to
 // update it every time the problem is modified, so we will compute it here for now
 template <typename i_t>
-i_t get_max_row_size(rmm::device_uvector<i_t>& offsets, rmm::cuda_stream_view stream_view)
+i_t get_max_row_size(rmm::device_uvector<i_t>& offsets, cuda::stream_ref stream_view)
 {
   auto begin = thrust::make_zip_iterator(thrust::make_tuple(offsets.begin(), offsets.begin() + 1));
   auto end   = thrust::make_zip_iterator(thrust::make_tuple(offsets.end() - 1, offsets.end()));
diff --git a/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cu b/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cu
index 290f6bef92..bf40838279 100644
--- a/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cu
+++ b/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cu
@@ -20,6 +20,7 @@
 #include <utilities/device_utils.cuh>
 
 #include <cub/cub.cuh>
+#include <cuda/stream>
 #include <raft/core/nvtx.hpp>
 #include "load_balanced_bounds_presolve.cuh"
 #include "load_balanced_bounds_presolve_helpers.cuh"
@@ -90,7 +91,7 @@ load_balanced_bounds_presolve_t<i_t, f_t>::~load_balanced_bounds_presolve_t()
 }
 
 template <typename i_t>
-std::pair<bool, i_t> sub_warp_meta(rmm::cuda_stream_view stream,
+std::pair<bool, i_t> sub_warp_meta(cuda::stream_ref stream,
                                    rmm::device_uvector<i_t>& d_warp_offsets,
                                    rmm::device_uvector<i_t>& d_warp_id_offsets,
                                    const std::vector<i_t>& bin_offsets,
diff --git a/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cuh b/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cuh
index 5e88f7bebf..b9266cfd97 100644
--- a/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cuh
+++ b/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cuh
@@ -7,6 +7,7 @@
 
 #pragma once
 
+#include <cuda/stream>
 #include "probing_cache.cuh"
 
 #include <mip_heuristics/problem/load_balanced_problem.cuh>
@@ -50,13 +51,13 @@ class managed_stream_pool {
   managed_stream_pool& operator=(managed_stream_pool const&) = delete;
 
   /**
-   * @brief Get a `cuda_stream_view` of a stream in the pool.
+   * @brief Get a `cuda::stream_ref` of a stream in the pool.
    *
    * This function is thread safe with respect to other calls to the same function.
    *
-   * @return rmm::cuda_stream_view
+   * @return cuda::stream_ref
    */
-  rmm::cuda_stream_view get_stream() const noexcept
+  cuda::stream_ref get_stream() const noexcept
   {
     int stream_id = (next_stream++) % streams_.size();
     end_unsycned  = std::max(stream_id, end_unsycned);
diff --git a/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve_helpers.cuh b/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve_helpers.cuh
index 7e3885b795..b36718b79d 100644
--- a/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve_helpers.cuh
+++ b/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve_helpers.cuh
@@ -15,7 +15,7 @@
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/scan.h>
-#include <rmm/cuda_stream_view.hpp>
+#include <cuda/stream>
 #include <rmm/device_uvector.hpp>
 #include <vector>
 
@@ -80,7 +80,7 @@ struct heavy_vertex_meta_t {
 };
 
 template <typename i_t>
-i_t create_heavy_item_block_segments(rmm::cuda_stream_view stream,
+i_t create_heavy_item_block_segments(cuda::stream_ref stream,
                                      rmm::device_uvector<i_t>& vertex_id,
                                      rmm::device_uvector<i_t>& pseudo_block_id,
                                      rmm::device_uvector<i_t>& item_block_segments,
diff --git a/cpp/src/mip_heuristics/presolve/semi_continuous.cu b/cpp/src/mip_heuristics/presolve/semi_continuous.cu
index 33b7efff0e..5ffac07b47 100644
--- a/cpp/src/mip_heuristics/presolve/semi_continuous.cu
+++ b/cpp/src/mip_heuristics/presolve/semi_continuous.cu
@@ -20,6 +20,7 @@
 #include <utilities/copy_helpers.hpp>
 #include <utilities/logger.hpp>
 
+#include <cuda/stream>
 #include <raft/util/cudart_utils.hpp>
 
 #include <algorithm>
@@ -372,7 +373,7 @@ template <typename i_t, typename f_t>
 void expand_initial_solutions_for_semi_continuous(
   mip_solver_settings_t<i_t, f_t>& settings,
   const std::vector<i_t>& semi_continuous_binary_to_original_indices,
-  rmm::cuda_stream_view stream)
+  cuda::stream_ref stream)
 {
   if (semi_continuous_binary_to_original_indices.empty()) { return; }
 
@@ -399,7 +400,7 @@ template void append_semi_continuous_auxiliaries_to_assignment(
 template void strip_semi_continuous_auxiliaries_from_assignment(std::vector<float>&, int);
 template void expand_initial_solutions_for_semi_continuous(mip_solver_settings_t<int, float>&,
                                                            const std::vector<int>&,
-                                                           rmm::cuda_stream_view);
+                                                           cuda::stream_ref);
 #endif
 
 #if MIP_INSTANTIATE_DOUBLE
@@ -412,7 +413,7 @@ template void append_semi_continuous_auxiliaries_to_assignment(
 template void strip_semi_continuous_auxiliaries_from_assignment(std::vector<double>&, int);
 template void expand_initial_solutions_for_semi_continuous(mip_solver_settings_t<int, double>&,
                                                            const std::vector<int>&,
-                                                           rmm::cuda_stream_view);
+                                                           cuda::stream_ref);
 #endif
 
 }  // namespace cuopt::mathematical_optimization::mip
diff --git a/cpp/src/mip_heuristics/presolve/semi_continuous.cuh b/cpp/src/mip_heuristics/presolve/semi_continuous.cuh
index ce3edc16a5..01540d8723 100644
--- a/cpp/src/mip_heuristics/presolve/semi_continuous.cuh
+++ b/cpp/src/mip_heuristics/presolve/semi_continuous.cuh
@@ -7,6 +7,7 @@
 
 #pragma once
 
+#include <cuda/stream>
 #include <cuopt/mathematical_optimization/mip/solver_settings.hpp>
 #include <cuopt/mathematical_optimization/optimization_problem.hpp>
 
@@ -55,7 +56,7 @@ template <typename i_t, typename f_t>
 void expand_initial_solutions_for_semi_continuous(
   mip_solver_settings_t<i_t, f_t>& settings,
   const std::vector<i_t>& semi_continuous_binary_to_original_indices,
-  rmm::cuda_stream_view stream);
+  cuda::stream_ref stream);
 
 template <typename i_t, typename f_t>
 void append_semi_continuous_auxiliaries_to_assignment(
diff --git a/cpp/src/mip_heuristics/presolve/third_party_presolve.cpp b/cpp/src/mip_heuristics/presolve/third_party_presolve.cpp
index 0292c8ff8f..61771100fa 100644
--- a/cpp/src/mip_heuristics/presolve/third_party_presolve.cpp
+++ b/cpp/src/mip_heuristics/presolve/third_party_presolve.cpp
@@ -48,6 +48,7 @@
 #include <utilities/macros.cuh>
 #include <utilities/timer.hpp>
 
+#include <cuda/stream>
 #include <raft/core/nvtx.hpp>
 
 #include <algorithm>
@@ -1207,7 +1208,7 @@ void third_party_presolve_t<i_t, f_t>::undo_from_device(rmm::device_uvector<f_t>
                                                         problem_category_t category,
                                                         bool status_to_skip,
                                                         bool dual_postsolve,
-                                                        rmm::cuda_stream_view stream_view)
+                                                        cuda::stream_ref stream_view)
 {
   std::vector<f_t> h_primal(primal_solution.size());
   std::vector<f_t> h_dual(dual_solution.size());
diff --git a/cpp/src/mip_heuristics/presolve/third_party_presolve.hpp b/cpp/src/mip_heuristics/presolve/third_party_presolve.hpp
index 79eb28d01e..f7b1ada5af 100644
--- a/cpp/src/mip_heuristics/presolve/third_party_presolve.hpp
+++ b/cpp/src/mip_heuristics/presolve/third_party_presolve.hpp
@@ -7,6 +7,7 @@
 
 #pragma once
 
+#include <cuda/stream>
 #include <memory>
 #include <optional>
 #include <string>
@@ -130,7 +131,7 @@ class third_party_presolve_t {
                         problem_category_t category,
                         bool status_to_skip,
                         bool dual_postsolve,
-                        rmm::cuda_stream_view stream_view);
+                        cuda::stream_ref stream_view);
 
   // Host-only postsolve. Resizes the vectors to original-problem dimensions.
   void undo(std::vector<f_t>& primal_solution,
diff --git a/cpp/src/mip_heuristics/problem/presolve_data.cu b/cpp/src/mip_heuristics/problem/presolve_data.cu
index 3c621bc2cd..6cad0afd20 100644
--- a/cpp/src/mip_heuristics/problem/presolve_data.cu
+++ b/cpp/src/mip_heuristics/problem/presolve_data.cu
@@ -17,6 +17,7 @@
 
 #include <mip_heuristics/utils.cuh>
 
+#include <cuda/stream>
 #include <raft/core/nvtx.hpp>
 
 #include <thrust/gather.h>
@@ -150,7 +151,7 @@ void presolve_data_t<i_t, f_t>::post_process_assignment(
   problem_t<i_t, f_t>& problem,
   rmm::device_uvector<f_t>& current_assignment,
   bool resize_to_original_problem,
-  rmm::cuda_stream_view stream)
+  cuda::stream_ref stream)
 {
   raft::common::nvtx::range fun_scope("post_process_assignment");
   cuopt_assert(current_assignment.size() == variable_mapping.size(), "size mismatch");
diff --git a/cpp/src/mip_heuristics/problem/presolve_data.cuh b/cpp/src/mip_heuristics/problem/presolve_data.cuh
index 713bb24c0d..9c112bdb13 100644
--- a/cpp/src/mip_heuristics/problem/presolve_data.cuh
+++ b/cpp/src/mip_heuristics/problem/presolve_data.cuh
@@ -11,6 +11,7 @@
 
 #include <thrust/sequence.h>
 #include <thrust/uninitialized_fill.h>
+#include <cuda/stream>
 #include <rmm/device_uvector.hpp>
 
 namespace cuopt {
@@ -53,7 +54,7 @@ struct postsolve_reconstruction_t {
 template <typename i_t, typename f_t>
 class presolve_data_t {
  public:
-  presolve_data_t(const optimization_problem_t<i_t, f_t>& problem, rmm::cuda_stream_view stream)
+  presolve_data_t(const optimization_problem_t<i_t, f_t>& problem, cuda::stream_ref stream)
     : variable_offsets(problem.get_n_variables(), 0),
       additional_var_used(problem.get_n_variables(), false),
       additional_var_id_per_var(problem.get_n_variables(), -1),
@@ -65,7 +66,7 @@ class presolve_data_t {
   {
   }
 
-  presolve_data_t(const presolve_data_t& other, rmm::cuda_stream_view stream)
+  presolve_data_t(const presolve_data_t& other, cuda::stream_ref stream)
     : variable_offsets(other.variable_offsets),
       additional_var_used(other.additional_var_used),
       additional_var_id_per_var(other.additional_var_id_per_var),
@@ -106,7 +107,7 @@ class presolve_data_t {
   void post_process_assignment(problem_t<i_t, f_t>& problem,
                                rmm::device_uvector<f_t>& current_assignment,
                                bool resize_to_original_problem,
-                               rmm::cuda_stream_view stream);
+                               cuda::stream_ref stream);
   void post_process_assignment(problem_t<i_t, f_t>& problem,
                                rmm::device_uvector<f_t>& current_assignment,
                                bool resize_to_original_problem = true)
diff --git a/cpp/src/mip_heuristics/problem/problem.cu b/cpp/src/mip_heuristics/problem/problem.cu
index 411724ede8..018a35f1f0 100644
--- a/cpp/src/mip_heuristics/problem/problem.cu
+++ b/cpp/src/mip_heuristics/problem/problem.cu
@@ -40,6 +40,7 @@
 #include <thrust/transform_reduce.h>
 #include <thrust/tuple.h>
 #include <cuda/std/functional>
+#include <cuda/stream>
 
 #include <raft/sparse/detail/cusparse_wrappers.h>
 #include <raft/core/logger.hpp>
@@ -2168,7 +2169,7 @@ bool problem_t<i_t, f_t>::pre_process_assignment(rmm::device_uvector<f_t>& assig
 template <typename i_t, typename f_t>
 void problem_t<i_t, f_t>::post_process_assignment(rmm::device_uvector<f_t>& current_assignment,
                                                   bool resize_to_original_problem,
-                                                  rmm::cuda_stream_view stream)
+                                                  cuda::stream_ref stream)
 {
   presolve_data.post_process_assignment(
     *this, current_assignment, resize_to_original_problem, stream);
diff --git a/cpp/src/mip_heuristics/problem/problem.cuh b/cpp/src/mip_heuristics/problem/problem.cuh
index 3ea3973d1d..7c2e27aa8c 100644
--- a/cpp/src/mip_heuristics/problem/problem.cuh
+++ b/cpp/src/mip_heuristics/problem/problem.cuh
@@ -26,6 +26,7 @@
 #include <utilities/macros.cuh>
 
 #include <branch_and_bound/constants.hpp>
+#include <cuda/stream>
 #include <memory>
 #include <raft/core/nvtx.hpp>
 #include <raft/random/rng_device.cuh>
@@ -102,7 +103,7 @@ class problem_t {
   bool pre_process_assignment(rmm::device_uvector<f_t>& assignment);
   void post_process_assignment(rmm::device_uvector<f_t>& current_assignment,
                                bool resize_to_original_problem,
-                               rmm::cuda_stream_view stream);
+                               cuda::stream_ref stream);
   void post_process_assignment(rmm::device_uvector<f_t>& current_assignment,
                                bool resize_to_original_problem = true)
   {
diff --git a/cpp/src/mip_heuristics/relaxed_lp/lp_state.cuh b/cpp/src/mip_heuristics/relaxed_lp/lp_state.cuh
index 6c9191ad3b..e655d15a84 100644
--- a/cpp/src/mip_heuristics/relaxed_lp/lp_state.cuh
+++ b/cpp/src/mip_heuristics/relaxed_lp/lp_state.cuh
@@ -8,6 +8,7 @@
 #pragma once
 
 #include <thrust/fill.h>
+#include <cuda/stream>
 #include <raft/util/cudart_utils.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -19,7 +20,7 @@ class problem_t;
 template <typename i_t, typename f_t>
 class lp_state_t {
  public:
-  lp_state_t(problem_t<i_t, f_t>& problem, rmm::cuda_stream_view stream)
+  lp_state_t(problem_t<i_t, f_t>& problem, cuda::stream_ref stream)
     : prev_primal(problem.n_variables, stream), prev_dual(problem.n_constraints, stream)
   {
     thrust::fill(
@@ -47,7 +48,7 @@ class lp_state_t {
   lp_state_t(lp_state_t<i_t, f_t>&& other) noexcept            = default;
   lp_state_t& operator=(lp_state_t<i_t, f_t>&& other) noexcept = default;
 
-  void resize(problem_t<i_t, f_t>& problem, rmm::cuda_stream_view stream)
+  void resize(problem_t<i_t, f_t>& problem, cuda::stream_ref stream)
   {
     prev_primal.resize(problem.n_variables, stream);
     prev_dual.resize(problem.n_constraints, stream);
diff --git a/cpp/src/mip_heuristics/solver_settings.cu b/cpp/src/mip_heuristics/solver_settings.cu
index a5325137bf..c25d876991 100644
--- a/cpp/src/mip_heuristics/solver_settings.cu
+++ b/cpp/src/mip_heuristics/solver_settings.cu
@@ -5,6 +5,7 @@
  */
 /* clang-format on */
 
+#include <cuda/stream>
 #include <cuopt/error.hpp>
 #include <cuopt/export.hpp>
 #include <cuopt/mathematical_optimization/mip/solver_settings.hpp>
@@ -16,7 +17,7 @@ namespace cuopt::mathematical_optimization {
 template <typename i_t, typename f_t>
 void mip_solver_settings_t<i_t, f_t>::add_initial_solution(const f_t* initial_solution,
                                                            i_t size,
-                                                           rmm::cuda_stream_view stream)
+                                                           cuda::stream_ref stream)
 {
   cuopt_expects(
     initial_solution != nullptr, error_type_t::ValidationError, "initial_solution cannot be null");
diff --git a/cpp/src/mip_heuristics/solver_solution.cu b/cpp/src/mip_heuristics/solver_solution.cu
index 8e89829fe3..f758d6b11d 100644
--- a/cpp/src/mip_heuristics/solver_solution.cu
+++ b/cpp/src/mip_heuristics/solver_solution.cu
@@ -10,6 +10,7 @@
 #include <mip_heuristics/mip_constants.hpp>
 #include <utilities/logger.hpp>
 
+#include <cuda/stream>
 #include <limits>
 #include <math_optimization/solution_writer.hpp>
 #include <raft/core/nvtx.hpp>
@@ -46,7 +47,7 @@ mip_solution_t<i_t, f_t>::mip_solution_t(rmm::device_uvector<f_t> solution,
 template <typename i_t, typename f_t>
 mip_solution_t<i_t, f_t>::mip_solution_t(mip_termination_status_t termination_status,
                                          solver_stats_t<i_t, f_t> stats,
-                                         rmm::cuda_stream_view stream_view)
+                                         cuda::stream_ref stream_view)
   : solution_(0, stream_view),
     objective_(0),
     mip_gap_(0),
@@ -61,7 +62,7 @@ mip_solution_t<i_t, f_t>::mip_solution_t(mip_termination_status_t termination_st
 
 template <typename i_t, typename f_t>
 mip_solution_t<i_t, f_t>::mip_solution_t(const cuopt::logic_error& error_status,
-                                         rmm::cuda_stream_view stream_view)
+                                         cuda::stream_ref stream_view)
   : solution_(0, stream_view),
     objective_(0),
     mip_gap_(0),
@@ -202,7 +203,7 @@ const std::vector<rmm::device_uvector<f_t>>& mip_solution_t<i_t, f_t>::get_solut
 
 template <typename i_t, typename f_t>
 void mip_solution_t<i_t, f_t>::write_to_sol_file(std::string_view filename,
-                                                 rmm::cuda_stream_view stream_view) const
+                                                 cuda::stream_ref stream_view) const
 {
   std::string status = get_termination_status_string();
   // Override for no termination
diff --git a/cpp/src/mip_heuristics/utils.cuh b/cpp/src/mip_heuristics/utils.cuh
index 76e9ce8a3d..3c0beb7111 100644
--- a/cpp/src/mip_heuristics/utils.cuh
+++ b/cpp/src/mip_heuristics/utils.cuh
@@ -10,6 +10,7 @@
 #include <thrust/inner_product.h>
 #include <thrust/logical.h>
 #include <thrust/transform_reduce.h>
+#include <cuda/stream>
 #include <cuopt/error.hpp>
 #include <pdlp/utils.cuh>
 #include <raft/random/rng_device.cuh>
@@ -31,7 +32,7 @@ constexpr int default_int_lower     = std::numeric_limits<int>::min();
 constexpr double zero_bound         = 0.;
 
 template <typename i_t>
-inline uint32_t compute_hash(raft::device_span<i_t> values, rmm::cuda_stream_view stream)
+inline uint32_t compute_hash(raft::device_span<i_t> values, cuda::stream_ref stream)
 {
   auto h_contents = cuopt::host_copy(values, stream);
   RAFT_CHECK_CUDA(stream.get());
@@ -39,7 +40,7 @@ inline uint32_t compute_hash(raft::device_span<i_t> values, rmm::cuda_stream_vie
 }
 
 template <typename i_t>
-inline uint32_t compute_hash(const rmm::device_uvector<i_t>& values, rmm::cuda_stream_view stream)
+inline uint32_t compute_hash(const rmm::device_uvector<i_t>& values, cuda::stream_ref stream)
 {
   auto h_contents = cuopt::host_copy(values, stream);
   RAFT_CHECK_CUDA(stream.get());
@@ -261,7 +262,7 @@ f_t compute_objective_from_vec(const rmm::device_uvector<f_t>& assignment,
 template <typename i_t, typename f_t>
 f_t compute_objective_from_vec(const rmm::device_uvector<f_t>& assignment,
                                const rmm::device_uvector<f_t>& objective_coefficients,
-                               rmm::cuda_stream_view stream)
+                               cuda::stream_ref stream)
 {
   cuopt_assert(assignment.size() == objective_coefficients.size(), "Size mismatch!");
   f_t computed_obj = thrust::inner_product(rmm::exec_policy(stream),
@@ -331,7 +332,7 @@ static __global__ void run_lambda_kernel(F f)
 // run a printf statement from the device side, useful for debugging without having to deal with
 // explicit memcpys
 template <typename Func>
-static void inline run_device_lambda(const rmm::cuda_stream_view& stream, Func f)
+static void inline run_device_lambda(const cuda::stream_ref& stream, Func f)
 {
   run_lambda_kernel<<<1, 1, 0, stream.get()>>>(f);
 }
diff --git a/cpp/src/pdlp/cpu_pdlp_warm_start_data.cu b/cpp/src/pdlp/cpu_pdlp_warm_start_data.cu
index 604c7b4377..8e9af5e3e2 100644
--- a/cpp/src/pdlp/cpu_pdlp_warm_start_data.cu
+++ b/cpp/src/pdlp/cpu_pdlp_warm_start_data.cu
@@ -5,6 +5,7 @@
  */
 /* clang-format on */
 
+#include <cuda/stream>
 #include <cuopt/export.hpp>
 #include <cuopt/mathematical_optimization/cpu_pdlp_warm_start_data.hpp>
 #include <cuopt/mathematical_optimization/pdlp/pdlp_warm_start_data.hpp>
@@ -16,7 +17,7 @@ namespace cuopt::mathematical_optimization {
 // Helper to copy device_uvector to std::vector (D2H)
 template <typename T>
 std::vector<T> device_to_host_vector(const rmm::device_uvector<T>& device_vec,
-                                     rmm::cuda_stream_view stream)
+                                     cuda::stream_ref stream)
 {
   if (device_vec.size() == 0) return std::vector<T>();
 
@@ -29,7 +30,7 @@ std::vector<T> device_to_host_vector(const rmm::device_uvector<T>& device_vec,
 // Helper to copy std::vector to device_uvector (H2D)
 template <typename T>
 rmm::device_uvector<T> host_to_device_vector(const std::vector<T>& host_vec,
-                                             rmm::cuda_stream_view stream)
+                                             cuda::stream_ref stream)
 {
   if (host_vec.empty()) return rmm::device_uvector<T>(0, stream);
 
@@ -42,7 +43,7 @@ rmm::device_uvector<T> host_to_device_vector(const std::vector<T>& host_vec,
 // Convert GPU → CPU warmstart (D2H copy)
 template <typename i_t, typename f_t>
 cpu_pdlp_warm_start_data_t<i_t, f_t> convert_to_cpu_warmstart(
-  const pdlp_warm_start_data_t<i_t, f_t>& gpu_data, rmm::cuda_stream_view stream)
+  const pdlp_warm_start_data_t<i_t, f_t>& gpu_data, cuda::stream_ref stream)
 {
   cpu_pdlp_warm_start_data_t<i_t, f_t> cpu_data;
 
@@ -77,7 +78,7 @@ cpu_pdlp_warm_start_data_t<i_t, f_t> convert_to_cpu_warmstart(
 // Convert CPU → GPU warmstart (H2D copy)
 template <typename i_t, typename f_t>
 pdlp_warm_start_data_t<i_t, f_t> convert_to_gpu_warmstart(
-  const cpu_pdlp_warm_start_data_t<i_t, f_t>& cpu_data, rmm::cuda_stream_view stream)
+  const cpu_pdlp_warm_start_data_t<i_t, f_t>& cpu_data, cuda::stream_ref stream)
 {
   pdlp_warm_start_data_t<i_t, f_t> gpu_data;
 
@@ -111,17 +112,17 @@ pdlp_warm_start_data_t<i_t, f_t> convert_to_gpu_warmstart(
 
 #if MIP_INSTANTIATE_DOUBLE
 template CUOPT_EXPORT cpu_pdlp_warm_start_data_t<int, double> convert_to_cpu_warmstart(
-  const pdlp_warm_start_data_t<int, double>&, rmm::cuda_stream_view);
+  const pdlp_warm_start_data_t<int, double>&, cuda::stream_ref);
 template CUOPT_EXPORT pdlp_warm_start_data_t<int, double> convert_to_gpu_warmstart(
-  const cpu_pdlp_warm_start_data_t<int, double>&, rmm::cuda_stream_view);
+  const cpu_pdlp_warm_start_data_t<int, double>&, cuda::stream_ref);
 #endif
 
 #if MIP_INSTANTIATE_FLOAT || PDLP_INSTANTIATE_FLOAT
 template CUOPT_EXPORT cpu_pdlp_warm_start_data_t<int, float> convert_to_cpu_warmstart(
-  const pdlp_warm_start_data_t<int, float>&, rmm::cuda_stream_view);
+  const pdlp_warm_start_data_t<int, float>&, cuda::stream_ref);
 
 template CUOPT_EXPORT pdlp_warm_start_data_t<int, float> convert_to_gpu_warmstart(
-  const cpu_pdlp_warm_start_data_t<int, float>&, rmm::cuda_stream_view);
+  const cpu_pdlp_warm_start_data_t<int, float>&, cuda::stream_ref);
 #endif
 
 }  // namespace cuopt::mathematical_optimization
diff --git a/cpp/src/pdlp/cuopt_c_internal.hpp b/cpp/src/pdlp/cuopt_c_internal.hpp
index 104dbc07af..6b41785819 100644
--- a/cpp/src/pdlp/cuopt_c_internal.hpp
+++ b/cpp/src/pdlp/cuopt_c_internal.hpp
@@ -16,11 +16,8 @@
 #include <cuopt/mathematical_optimization/pdlp/solver_solution.hpp>
 
 #include <cuda/stream>
-
 #include <raft/core/handle.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
-
 #include <memory>
 
 namespace cuopt::mathematical_optimization {
@@ -31,8 +28,7 @@ struct problem_and_stream_view_t {
   {
     if (mem_backend == memory_backend_t::GPU) {
       // Use RAII locals so partial allocations are cleaned up if a later new throws
-      std::unique_ptr<rmm::cuda_stream_view> sv(
-        new rmm::cuda_stream_view(cuda::stream_ref{cudaStreamPerThread}));
+      std::unique_ptr<cuda::stream_ref> sv(new cuda::stream_ref{cudaStreamPerThread});
       std::unique_ptr<raft::handle_t> h(new raft::handle_t(*sv));
       std::unique_ptr<optimization_problem_t<cuopt_int_t, cuopt_float_t>> gp(
         new optimization_problem_t<cuopt_int_t, cuopt_float_t>(h.get()));
@@ -117,9 +113,8 @@ struct problem_and_stream_view_t {
   memory_backend_t memory_backend;
   optimization_problem_t<cuopt_int_t, cuopt_float_t>* gpu_problem;
   cpu_optimization_problem_t<cuopt_int_t, cuopt_float_t>* cpu_problem;
-  rmm::cuda_stream_view*
-    stream_view_ptr;           // nullptr for CPU memory backend to avoid CUDA initialization
-  raft::handle_t* handle_ptr;  // nullptr for CPU memory backend to avoid CUDA initialization
+  cuda::stream_ref* stream_view_ptr;  // nullptr for CPU memory backend to avoid CUDA initialization
+  raft::handle_t* handle_ptr;         // nullptr for CPU memory backend to avoid CUDA initialization
 };
 
 struct solution_and_stream_view_t {
diff --git a/cpp/src/pdlp/distributed_pdlp/multi_gpu_engine.cu b/cpp/src/pdlp/distributed_pdlp/multi_gpu_engine.cu
index 70f32bc5b0..40cfa4855d 100644
--- a/cpp/src/pdlp/distributed_pdlp/multi_gpu_engine.cu
+++ b/cpp/src/pdlp/distributed_pdlp/multi_gpu_engine.cu
@@ -8,6 +8,7 @@
 
 #include <cuopt/error.hpp>
 
+#include <cuda/stream>
 #include <raft/core/device_setter.hpp>
 
 #include <nccl.h>
@@ -111,15 +112,14 @@ void multi_gpu_engine_t<i_t, f_t>::synchronize_shards()
 }
 
 template <typename i_t, typename f_t>
-void multi_gpu_engine_t<i_t, f_t>::graph_capture_fork_to_shards(rmm::cuda_stream_view master_stream)
+void multi_gpu_engine_t<i_t, f_t>::graph_capture_fork_to_shards(cuda::stream_ref master_stream)
 {
   graph_master_ready_event_->record(master_stream);
   for_each_shard([&](auto& s) { graph_master_ready_event_->stream_wait(s.stream.view()); });
 }
 
 template <typename i_t, typename f_t>
-void multi_gpu_engine_t<i_t, f_t>::graph_capture_join_from_shards(
-  rmm::cuda_stream_view master_stream)
+void multi_gpu_engine_t<i_t, f_t>::graph_capture_join_from_shards(cuda::stream_ref master_stream)
 {
   for_each_shard([&](auto& s, int r) { graph_shard_ready_events_[r]->record(s.stream.view()); });
   for (auto& e : graph_shard_ready_events_) {
@@ -128,14 +128,14 @@ void multi_gpu_engine_t<i_t, f_t>::graph_capture_join_from_shards(
 }
 
 template <typename i_t, typename f_t>
-void multi_gpu_engine_t<i_t, f_t>::sync_await_master(rmm::cuda_stream_view master_stream)
+void multi_gpu_engine_t<i_t, f_t>::sync_await_master(cuda::stream_ref master_stream)
 {
   sync_master_ready_event_->record(master_stream);
   for_each_shard([&](auto& s) { sync_master_ready_event_->stream_wait(s.stream.view()); });
 }
 
 template <typename i_t, typename f_t>
-void multi_gpu_engine_t<i_t, f_t>::sync_await_shards(rmm::cuda_stream_view master_stream)
+void multi_gpu_engine_t<i_t, f_t>::sync_await_shards(cuda::stream_ref master_stream)
 {
   for_each_shard([&](auto& s, int r) { sync_shard_ready_events_[r]->record(s.stream.view()); });
   for (auto& e : sync_shard_ready_events_) {
diff --git a/cpp/src/pdlp/distributed_pdlp/multi_gpu_engine.hpp b/cpp/src/pdlp/distributed_pdlp/multi_gpu_engine.hpp
index d3896f4ce7..da22ee556e 100644
--- a/cpp/src/pdlp/distributed_pdlp/multi_gpu_engine.hpp
+++ b/cpp/src/pdlp/distributed_pdlp/multi_gpu_engine.hpp
@@ -15,6 +15,7 @@
 #include <cuopt/mathematical_optimization/pdlp/solver_settings.hpp>
 
 #include <raft/sparse/detail/cusparse_wrappers.h>
+#include <cuda/stream>
 #include <raft/core/cusparse_macros.hpp>
 #include <raft/core/device_mdspan.hpp>
 #include <raft/core/device_setter.hpp>
@@ -536,18 +537,18 @@ struct multi_gpu_engine_t {
   std::vector<std::unique_ptr<cuopt::event_handler_t>> sync_shard_ready_events_;
 
   // Forks master stream to shards, so that the captured graph can see the work on the shards
-  void graph_capture_fork_to_shards(rmm::cuda_stream_view master_stream);
+  void graph_capture_fork_to_shards(cuda::stream_ref master_stream);
 
   // Joins shards back to master stream for correct graph capture
-  void graph_capture_join_from_shards(rmm::cuda_stream_view master_stream);
+  void graph_capture_join_from_shards(cuda::stream_ref master_stream);
 
   // Functionnaly same as graph_capture_fork_to_shards but on a different event to avoid race
   // conditions Can be used as a way to sync shards with master stream
-  void sync_await_master(rmm::cuda_stream_view master_stream);
+  void sync_await_master(cuda::stream_ref master_stream);
 
   // Same as sync_await_master
   // Can be used as a way to sync master stream with shards
-  void sync_await_shards(rmm::cuda_stream_view master_stream);
+  void sync_await_shards(cuda::stream_ref master_stream);
 };
 
 }  // namespace cuopt::mathematical_optimization::pdlp
diff --git a/cpp/src/pdlp/initial_scaling_strategy/initial_scaling.cuh b/cpp/src/pdlp/initial_scaling_strategy/initial_scaling.cuh
index df67024630..96d7f0629c 100644
--- a/cpp/src/pdlp/initial_scaling_strategy/initial_scaling.cuh
+++ b/cpp/src/pdlp/initial_scaling_strategy/initial_scaling.cuh
@@ -13,9 +13,9 @@
 
 #include <mip_heuristics/solution/solution.cuh>
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 
 #include <limits>
@@ -146,7 +146,7 @@ class pdlp_initial_scaling_strategy_t {
   void reset_integer_variables();
 
   raft::handle_t const* handle_ptr_{nullptr};
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
 
   i_t primal_size_h_;
   i_t dual_size_h_;
diff --git a/cpp/src/pdlp/optimal_batch_size_handler/optimal_batch_size_handler.cu b/cpp/src/pdlp/optimal_batch_size_handler/optimal_batch_size_handler.cu
index 0df699ae1e..19cc046f15 100644
--- a/cpp/src/pdlp/optimal_batch_size_handler/optimal_batch_size_handler.cu
+++ b/cpp/src/pdlp/optimal_batch_size_handler/optimal_batch_size_handler.cu
@@ -12,6 +12,7 @@
 #include <utilities/device_scalar_init.hpp>
 #include <utilities/event_handler.cuh>
 
+#include <cuda/stream>
 #include <raft/core/cusparse_macros.hpp>
 
 #include <mip_heuristics/mip_constants.hpp>
@@ -169,7 +170,7 @@ static double evaluate_node(cusparse_sp_mat_descr_view A,
 {
   cuopt_assert(current_batch_size > 0, "Current batch size must be greater than 0");
 
-  rmm::cuda_stream_view stream_view = handle_ptr->get_stream();
+  cuda::stream_ref stream_view = handle_ptr->get_stream();
   SpMM_benchmarks_context_t<i_t, f_t> spmm_benchmarks_context(
     A, A_T, primal_size, dual_size, current_batch_size, handle_ptr);
 
@@ -219,7 +220,7 @@ int optimal_batch_size_handler(const optimization_problem_t<i_t, f_t>& op_proble
     std::pow(2, std::floor(std::log2(std::min(initial_batch_size, max_batch_size))));
   int optimal_batch_size = current_batch_size;
   double best_ratio;
-  rmm::cuda_stream_view stream_view = op_problem.get_handle_ptr()->get_stream();
+  cuda::stream_ref stream_view = op_problem.get_handle_ptr()->get_stream();
 
   mip::problem_t<i_t, f_t> problem(op_problem);
 
diff --git a/cpp/src/pdlp/optimization_problem.cu b/cpp/src/pdlp/optimization_problem.cu
index 87fe438ca4..e6c77e383f 100644
--- a/cpp/src/pdlp/optimization_problem.cu
+++ b/cpp/src/pdlp/optimization_problem.cu
@@ -21,6 +21,7 @@
 #include <utilities/logger.hpp>
 #include <utilities/sparse_matrix_helpers.hpp>
 
+#include <cuda/stream>
 #include <raft/core/copy.hpp>
 #include <raft/core/cuda_support.hpp>
 #include <raft/core/device_mdspan.hpp>
@@ -29,7 +30,6 @@
 #include <raft/util/cuda_utils.cuh>
 #include <raft/util/cudart_utils.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
@@ -57,7 +57,7 @@ namespace cuopt::mathematical_optimization {
 template <typename i_t, typename f_t>
 optimization_problem_t<i_t, f_t>::optimization_problem_t(raft::handle_t const* handle_ptr)
   : handle_ptr_(handle_ptr),
-    stream_view_(handle_ptr != nullptr ? handle_ptr->get_stream() : rmm::cuda_stream_view{}),
+    stream_view_(handle_ptr != nullptr ? handle_ptr->get_stream() : cuda::stream_ref{}),
     A_(0, stream_view_),
     A_indices_(0, stream_view_),
     A_offsets_(0, stream_view_),
@@ -1015,7 +1015,7 @@ static bool csr_matrices_equivalent_with_permutation(const rmm::device_uvector<i
                                                      const rmm::device_uvector<i_t>& d_row_perm_inv,
                                                      const rmm::device_uvector<i_t>& d_col_perm_inv,
                                                      i_t n_cols,
-                                                     rmm::cuda_stream_view stream)
+                                                     cuda::stream_ref stream)
 {
   const i_t nnz = static_cast<i_t>(this_values.size());
   if (nnz != static_cast<i_t>(other_values.size())) { return false; }
@@ -1530,7 +1530,7 @@ struct cast_op {
 };
 
 template <typename From, typename To>
-rmm::device_uvector<To> gpu_cast(const rmm::device_uvector<From>& src, rmm::cuda_stream_view stream)
+rmm::device_uvector<To> gpu_cast(const rmm::device_uvector<From>& src, cuda::stream_ref stream)
 {
   rmm::device_uvector<To> dst(src.size(), stream);
   if (src.size() > 0) {
@@ -1541,14 +1541,14 @@ rmm::device_uvector<To> gpu_cast(const rmm::device_uvector<From>& src, rmm::cuda
 }
 
 template rmm::device_uvector<float> gpu_cast<double, float>(const rmm::device_uvector<double>&,
-                                                            rmm::cuda_stream_view);
+                                                            cuda::stream_ref);
 template rmm::device_uvector<double> gpu_cast<float, double>(const rmm::device_uvector<float>&,
-                                                             rmm::cuda_stream_view);
+                                                             cuda::stream_ref);
 
 template <typename i_t, typename f_t>
 template <typename other_f_t>
 optimization_problem_t<i_t, other_f_t> optimization_problem_t<i_t, f_t>::convert_to_other_prec(
-  rmm::cuda_stream_view stream) const
+  cuda::stream_ref stream) const
 {
   optimization_problem_t<i_t, other_f_t> other(handle_ptr_);
 
@@ -1633,8 +1633,7 @@ template class CUOPT_EXPORT optimization_problem_t<int32_t, double>;
 
 #if PDLP_INSTANTIATE_FLOAT || MIP_INSTANTIATE_FLOAT
 template CUOPT_EXPORT optimization_problem_t<int32_t, float>
-  optimization_problem_t<int32_t, double>::convert_to_other_prec<float>(
-    rmm::cuda_stream_view) const;
+  optimization_problem_t<int32_t, double>::convert_to_other_prec<float>(cuda::stream_ref) const;
 #endif
 
 }  // namespace cuopt::mathematical_optimization
diff --git a/cpp/src/pdlp/pdhg.cu b/cpp/src/pdlp/pdhg.cu
index 7ef0a615c0..586444679d 100644
--- a/cpp/src/pdlp/pdhg.cu
+++ b/cpp/src/pdlp/pdhg.cu
@@ -13,6 +13,7 @@
 #include <pdlp/utilities/ping_pong_graph.cuh>
 #include <pdlp/utils.cuh>
 
+#include <cuda/stream>
 #include <raft/core/device_span.hpp>
 
 #include <mip_heuristics/mip_constants.hpp>
@@ -172,7 +173,7 @@ new_bounds_groups_t<i_t, f_t> copy_new_bounds_to_groups(
   const rmm::device_uvector<f_t>& new_bounds_lower,
   const rmm::device_uvector<f_t>& new_bounds_upper,
   i_t batch_size,
-  rmm::cuda_stream_view stream_view)
+  cuda::stream_ref stream_view)
 {
   cuopt_assert(new_bounds_climber_id.size() == new_bounds_idx.size(),
                "New bounds climber id and index sizes must match");
@@ -210,7 +211,7 @@ void copy_groups_to_new_bounds(const new_bounds_groups_t<i_t, f_t>& groups,
                                rmm::device_uvector<i_t>& new_bounds_idx,
                                rmm::device_uvector<f_t>& new_bounds_lower,
                                rmm::device_uvector<f_t>& new_bounds_upper,
-                               rmm::cuda_stream_view stream_view)
+                               cuda::stream_ref stream_view)
 {
   size_t n_entries = 0;
   for (i_t c = 0; c < group_count; ++c) {
diff --git a/cpp/src/pdlp/pdhg.hpp b/cpp/src/pdlp/pdhg.hpp
index f6eb9931fa..fdc8b82dfe 100644
--- a/cpp/src/pdlp/pdhg.hpp
+++ b/cpp/src/pdlp/pdhg.hpp
@@ -14,9 +14,9 @@
 #include <pdlp/swap_and_resize_helper.cuh>
 #include <pdlp/utilities/ping_pong_graph.cuh>
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -138,7 +138,7 @@ class pdhg_solver_t {
 
   bool batch_mode_{false};
   raft::handle_t const* handle_ptr_{nullptr};
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
 
   mip::problem_t<i_t, f_t>* problem_ptr;
 
diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu
index 86a32f028e..9f3365d59c 100644
--- a/cpp/src/pdlp/pdlp.cu
+++ b/cpp/src/pdlp/pdlp.cu
@@ -27,6 +27,7 @@
 #include <utilities/macros.cuh>
 
 #include <raft/sparse/detail/cusparse_wrappers.h>
+#include <cuda/stream>
 #include <raft/core/cusparse_macros.hpp>
 #include <raft/core/nvtx.hpp>
 #include <raft/linalg/eltwise.cuh>
@@ -2187,7 +2188,7 @@ void pdlp_solver_t<i_t, f_t>::resize_and_swap_all_context_loop(
 // saddle-point delta buffers. Shared by the single-GPU and per-shard
 // (distributed) paths so the two only differ by which pdhg/stream they pass.
 template <typename i_t, typename f_t>
-static void compute_primal_dual_deltas(pdhg_solver_t<i_t, f_t>& pdhg, rmm::cuda_stream_view stream)
+static void compute_primal_dual_deltas(pdhg_solver_t<i_t, f_t>& pdhg, cuda::stream_ref stream)
 {
   cub::DeviceTransform::Transform(
     cuda::std::make_tuple(pdhg.get_reflected_primal().data(), pdhg.get_primal_solution().data()),
diff --git a/cpp/src/pdlp/pdlp.cuh b/cpp/src/pdlp/pdlp.cuh
index 60b2c1c354..74179a64a0 100644
--- a/cpp/src/pdlp/pdlp.cuh
+++ b/cpp/src/pdlp/pdlp.cuh
@@ -27,6 +27,7 @@
 
 #include <utilities/timer.hpp>
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
 #include <rmm/device_scalar.hpp>
@@ -176,7 +177,7 @@ class pdlp_solver_t {
   bool batch_mode_{false};
 
   raft::handle_t const* handle_ptr_;
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
   // Intentionnaly take a copy to avoid an unintentional modification in the calling context
   const pdlp_solver_settings_t<i_t, f_t> settings_;
   mip::shared_strong_branching_context_view_t<i_t, f_t> sb_view_{settings_.shared_sb_solved};
diff --git a/cpp/src/pdlp/pdlp_warm_start_data.cu b/cpp/src/pdlp/pdlp_warm_start_data.cu
index 2ce6b0f4b5..711a29c606 100644
--- a/cpp/src/pdlp/pdlp_warm_start_data.cu
+++ b/cpp/src/pdlp/pdlp_warm_start_data.cu
@@ -12,7 +12,6 @@
 #include <mip_heuristics/mip_constants.hpp>
 
 #include <cuda/stream>
-
 #include <rmm/device_uvector.hpp>
 
 #include <raft/util/cudart_utils.hpp>
@@ -88,7 +87,7 @@ pdlp_warm_start_data_t<i_t, f_t>::pdlp_warm_start_data_t()
 
 template <typename i_t, typename f_t>
 pdlp_warm_start_data_t<i_t, f_t>::pdlp_warm_start_data_t(
-  const pdlp_warm_start_data_view_t<i_t, f_t>& other, rmm::cuda_stream_view stream_view)
+  const pdlp_warm_start_data_view_t<i_t, f_t>& other, cuda::stream_ref stream_view)
   : current_primal_solution_(other.current_primal_solution_.size(), stream_view),
     current_dual_solution_(other.current_dual_solution_.size(), stream_view),
     initial_primal_average_(other.initial_primal_average_.size(), stream_view),
diff --git a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu
index 06158f4d1c..dc58066fc5 100644
--- a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu
+++ b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu
@@ -24,6 +24,7 @@
 #endif
 
 #include <raft/sparse/detail/cusparse_wrappers.h>
+#include <cuda/stream>
 #include <raft/core/device_span.hpp>
 #include <raft/core/nvtx.hpp>
 #include <raft/linalg/binary_op.cuh>
@@ -981,7 +982,7 @@ void pdlp_restart_strategy_t<i_t, f_t>::cupdlpx_restart(
   // Small copy helper to use in both single-GPU and distributed paths.
   auto commit_potential_next_as_last_restart = [](pdlp_restart_strategy_t<i_t, f_t>& rest,
                                                   pdhg_solver_t<i_t, f_t>& solver,
-                                                  rmm::cuda_stream_view stream) {
+                                                  cuda::stream_ref stream) {
     raft::copy(rest.last_restart_duality_gap_.primal_solution_.data(),
                solver.get_potential_next_primal_solution().data(),
                rest.last_restart_duality_gap_.primal_solution_.size(),
diff --git a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cuh b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cuh
index 68ef7503a0..f70358f550 100644
--- a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cuh
+++ b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cuh
@@ -19,9 +19,9 @@
 
 #include <mip_heuristics/problem/problem.cuh>
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_buffer.hpp>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
@@ -306,7 +306,7 @@ class pdlp_restart_strategy_t {
                                  rmm::device_uvector<f_t>& dual_step_size);
 
   raft::handle_t const* handle_ptr_{nullptr};
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
 
  public:
   const bool batch_mode_{false};
diff --git a/cpp/src/pdlp/restart_strategy/weighted_average_solution.hpp b/cpp/src/pdlp/restart_strategy/weighted_average_solution.hpp
index 777f9ea804..b9f9ab53cc 100644
--- a/cpp/src/pdlp/restart_strategy/weighted_average_solution.hpp
+++ b/cpp/src/pdlp/restart_strategy/weighted_average_solution.hpp
@@ -9,9 +9,9 @@
 #include <pdlp/saddle_point.hpp>
 #include <pdlp/utilities/ping_pong_graph.cuh>
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -36,7 +36,7 @@ class weighted_average_solution_t {
 
  private:
   raft::handle_t const* handle_ptr_{nullptr};
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
 
   i_t primal_size_h_;
   i_t dual_size_h_;
diff --git a/cpp/src/pdlp/saddle_point.cu b/cpp/src/pdlp/saddle_point.cu
index 5edf3a5c66..ef09ce095d 100644
--- a/cpp/src/pdlp/saddle_point.cu
+++ b/cpp/src/pdlp/saddle_point.cu
@@ -5,6 +5,7 @@
  */
 /* clang-format on */
 
+#include <cuda/stream>
 #include <cuopt/error.hpp>
 
 #include <pdlp/restart_strategy/pdlp_restart_strategy.cuh>
@@ -99,7 +100,7 @@ void saddle_point_state_t<i_t, f_t>::resize_context(i_t new_size)
 
 template <typename i_t, typename f_t>
 void saddle_point_state_t<i_t, f_t>::copy(saddle_point_state_t<i_t, f_t>& other,
-                                          rmm::cuda_stream_view stream)
+                                          cuda::stream_ref stream)
 {
   EXE_CUOPT_EXPECTS(this->primal_size_ == other.get_primal_size(),
                     "Size of primal solution must be the same in order to copy");
diff --git a/cpp/src/pdlp/saddle_point.hpp b/cpp/src/pdlp/saddle_point.hpp
index c6fc962016..55ac5a87d4 100644
--- a/cpp/src/pdlp/saddle_point.hpp
+++ b/cpp/src/pdlp/saddle_point.hpp
@@ -9,9 +9,9 @@
 
 #include <cuopt/mathematical_optimization/pdlp/pdlp_hyper_params.cuh>
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -81,7 +81,7 @@ class saddle_point_state_t {
    *
    * @throws cuopt::logic_error if the solutions are not of the same size
    */
-  void copy(saddle_point_state_t<i_t, f_t>& other, rmm::cuda_stream_view stream);
+  void copy(saddle_point_state_t<i_t, f_t>& other, cuda::stream_ref stream);
 
   i_t get_primal_size() const;
   i_t get_dual_size() const;
diff --git a/cpp/src/pdlp/solve.cu b/cpp/src/pdlp/solve.cu
index b995dc3f12..40b8d80378 100644
--- a/cpp/src/pdlp/solve.cu
+++ b/cpp/src/pdlp/solve.cu
@@ -49,7 +49,6 @@
 #include <pdlp/utilities/problem_checking.cuh>
 
 #include <raft/sparse/detail/cusparse_wrappers.h>
-
 #include <cuda/stream>
 #include <raft/core/cusparse_macros.hpp>
 #include <raft/core/device_setter.hpp>
@@ -76,7 +75,7 @@ namespace cuopt::mathematical_optimization {
 
 template <typename From, typename To>
 extern rmm::device_uvector<To> gpu_cast(const rmm::device_uvector<From>& src,
-                                        rmm::cuda_stream_view stream);
+                                        cuda::stream_ref stream);
 
 // This serves as both a warm up but also a mandatory initial call to setup cuSparse and cuBLAS
 static void init_handler(const raft::handle_t* handle_ptr)
@@ -338,7 +337,7 @@ std::atomic<int> global_concurrent_halt{0};
 template <typename f_t>
 void adjust_dual_solution_and_reduced_cost(rmm::device_uvector<f_t>& dual_solution,
                                            rmm::device_uvector<f_t>& reduced_cost,
-                                           rmm::cuda_stream_view stream_view)
+                                           cuda::stream_ref stream_view)
 {
   // y <- -y
   cub::DeviceTransform::Transform(
@@ -1283,9 +1282,9 @@ template <typename i_t, typename f_t>
 static optimization_problem_solution_t<i_t, f_t> run_batch_pdlp_splitting(
   optimization_problem_t<i_t, f_t>& problem, pdlp_solver_settings_t<i_t, f_t> const& settings)
 {
-  rmm::cuda_stream_view stream = problem.get_handle_ptr()->get_stream();
-  const i_t n_vars             = problem.get_n_variables();
-  const i_t n_constraints      = problem.get_n_constraints();
+  cuda::stream_ref stream = problem.get_handle_ptr()->get_stream();
+  const i_t n_vars        = problem.get_n_variables();
+  const i_t n_constraints = problem.get_n_constraints();
 
   // Splitting path only supports un-expanded problems + per-climber variable-bound overrides.
   cuopt_expects(problem.get_objective_coefficients().size() == static_cast<size_t>(n_vars),
@@ -1610,8 +1609,8 @@ optimization_problem_solution_t<i_t, f_t> run_concurrent(
         {
           try {
             auto call_barrier_thread = [&]() {
-              rmm::cuda_stream_view barrier_stream = cuda::stream_ref{cudaStreamPerThread};
-              barrier_handle_ptr = std::make_unique<raft::handle_t>(barrier_stream);
+              cuda::stream_ref barrier_stream = cuda::stream_ref{cudaStreamPerThread};
+              barrier_handle_ptr              = std::make_unique<raft::handle_t>(barrier_stream);
               run_barrier_thread<i_t, f_t>(dual_simplex_problem,
                                            settings_pdlp,
                                            sol_barrier_ptr,
@@ -1932,7 +1931,7 @@ optimization_problem_solution_t<i_t, f_t> solve_qcqp(
 template <typename i_t, typename f_t>
 static std::optional<optimization_problem_solution_t<i_t, f_t>>
 terminal_solution_from_presolve_status(mip::third_party_presolve_status_t status,
-                                       rmm::cuda_stream_view stream)
+                                       cuda::stream_ref stream)
 {
   switch (status) {
     case mip::third_party_presolve_status_t::INFEASIBLE:
diff --git a/cpp/src/pdlp/solver_settings.cu b/cpp/src/pdlp/solver_settings.cu
index 33d8f1a64b..e9f1813fdf 100644
--- a/cpp/src/pdlp/solver_settings.cu
+++ b/cpp/src/pdlp/solver_settings.cu
@@ -13,6 +13,7 @@
 #include <mip_heuristics/mip_constants.hpp>
 #include <utilities/logger.hpp>
 
+#include <cuda/stream>
 #include <raft/util/cudart_utils.hpp>
 
 #include <thrust/scatter.h>
@@ -36,7 +37,7 @@ void pdlp_solver_settings_t<i_t, f_t>::set_optimality_tolerance(f_t eps_optimal)
 
 template <typename i_t, typename f_t>
 void pdlp_solver_settings_t<i_t, f_t>::set_initial_primal_solution(
-  const f_t* initial_primal_solution, i_t size, rmm::cuda_stream_view stream)
+  const f_t* initial_primal_solution, i_t size, cuda::stream_ref stream)
 {
   cuopt_expects(initial_primal_solution != nullptr,
                 error_type_t::ValidationError,
@@ -49,7 +50,7 @@ void pdlp_solver_settings_t<i_t, f_t>::set_initial_primal_solution(
 template <typename i_t, typename f_t>
 void pdlp_solver_settings_t<i_t, f_t>::set_initial_dual_solution(const f_t* initial_dual_solution,
                                                                  i_t size,
-                                                                 rmm::cuda_stream_view stream)
+                                                                 cuda::stream_ref stream)
 {
   cuopt_expects(initial_dual_solution != nullptr,
                 error_type_t::ValidationError,
diff --git a/cpp/src/pdlp/solver_solution.cu b/cpp/src/pdlp/solver_solution.cu
index 0fbda1701e..ace97647f1 100644
--- a/cpp/src/pdlp/solver_solution.cu
+++ b/cpp/src/pdlp/solver_solution.cu
@@ -15,6 +15,7 @@
 #include <utilities/logger.hpp>
 #include <utilities/macros.cuh>
 
+#include <cuda/stream>
 #include <raft/core/nvtx.hpp>
 #include <raft/util/cudart_utils.hpp>
 
@@ -25,7 +26,7 @@ namespace cuopt::mathematical_optimization {
 
 template <typename i_t, typename f_t>
 optimization_problem_solution_t<i_t, f_t>::optimization_problem_solution_t(
-  pdlp_termination_status_t termination_status, rmm::cuda_stream_view stream_view)
+  pdlp_termination_status_t termination_status, cuda::stream_ref stream_view)
   : primal_solution_{0, stream_view},
     dual_solution_{0, stream_view},
     reduced_cost_{0, stream_view},
@@ -38,7 +39,7 @@ optimization_problem_solution_t<i_t, f_t>::optimization_problem_solution_t(
 
 template <typename i_t, typename f_t>
 optimization_problem_solution_t<i_t, f_t>::optimization_problem_solution_t(
-  cuopt::logic_error error_status_, rmm::cuda_stream_view stream_view)
+  cuopt::logic_error error_status_, cuda::stream_ref stream_view)
   : primal_solution_{0, stream_view},
     dual_solution_{0, stream_view},
     reduced_cost_{0, stream_view},
@@ -210,7 +211,7 @@ void optimization_problem_solution_t<i_t, f_t>::write_additional_termination_sta
 
 template <typename i_t, typename f_t>
 void optimization_problem_solution_t<i_t, f_t>::write_to_file(std::string_view filename,
-                                                              rmm::cuda_stream_view stream_view,
+                                                              cuda::stream_ref stream_view,
                                                               bool generate_variable_values)
 {
   raft::common::nvtx::range fun_scope("write final solution to file");
@@ -430,7 +431,7 @@ optimization_problem_solution_t<i_t, f_t>::get_pdlp_warm_start_data()
 
 template <typename i_t, typename f_t>
 void optimization_problem_solution_t<i_t, f_t>::write_to_sol_file(
-  std::string_view filename, rmm::cuda_stream_view stream_view) const
+  std::string_view filename, cuda::stream_ref stream_view) const
 {
   cuopt_expects(termination_stats_.size() == 1,
                 error_type_t::ValidationError,
diff --git a/cpp/src/pdlp/step_size_strategy/adaptive_step_size_strategy.hpp b/cpp/src/pdlp/step_size_strategy/adaptive_step_size_strategy.hpp
index ebe546c2de..5b777c9f74 100644
--- a/cpp/src/pdlp/step_size_strategy/adaptive_step_size_strategy.hpp
+++ b/cpp/src/pdlp/step_size_strategy/adaptive_step_size_strategy.hpp
@@ -16,9 +16,9 @@
 #include <pdlp/saddle_point.hpp>
 #include <pdlp/swap_and_resize_helper.cuh>
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -106,7 +106,7 @@ class adaptive_step_size_strategy_t {
   const bool batch_mode_;
 
   raft::handle_t const* handle_ptr_{nullptr};
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
 
   i_t primal_size_;
   i_t dual_size_;
diff --git a/cpp/src/pdlp/termination_strategy/convergence_information.hpp b/cpp/src/pdlp/termination_strategy/convergence_information.hpp
index 1bcb2fc0ab..4cd0307fca 100644
--- a/cpp/src/pdlp/termination_strategy/convergence_information.hpp
+++ b/cpp/src/pdlp/termination_strategy/convergence_information.hpp
@@ -18,9 +18,9 @@
 
 #include <mip_heuristics/problem/problem.cuh>
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_buffer.hpp>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
@@ -190,7 +190,7 @@ class convergence_information_t {
 
   raft::handle_t const* handle_ptr_{nullptr};
 
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
 
   i_t primal_size_h_;
   i_t dual_size_h_;
diff --git a/cpp/src/pdlp/termination_strategy/infeasibility_information.hpp b/cpp/src/pdlp/termination_strategy/infeasibility_information.hpp
index 2ccae0633f..2061080361 100644
--- a/cpp/src/pdlp/termination_strategy/infeasibility_information.hpp
+++ b/cpp/src/pdlp/termination_strategy/infeasibility_information.hpp
@@ -16,9 +16,9 @@
 
 #include <mip_heuristics/problem/problem.cuh>
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_buffer.hpp>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
@@ -85,7 +85,7 @@ class infeasibility_information_t {
   void compute_reduced_costs_dual_objective_contribution();
 
   raft::handle_t const* handle_ptr_{nullptr};
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
 
   i_t primal_size_h_;
   i_t dual_size_h_;
diff --git a/cpp/src/pdlp/termination_strategy/termination_strategy.hpp b/cpp/src/pdlp/termination_strategy/termination_strategy.hpp
index 948f7674d0..2c26303f50 100644
--- a/cpp/src/pdlp/termination_strategy/termination_strategy.hpp
+++ b/cpp/src/pdlp/termination_strategy/termination_strategy.hpp
@@ -19,9 +19,9 @@
 
 #include <utilities/unique_pinned_ptr.hpp>
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -214,7 +214,7 @@ class pdlp_termination_strategy_t {
   void check_termination_criteria();
 
   raft::handle_t const* handle_ptr_{nullptr};
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
 
   mip::problem_t<i_t, f_t>* problem_ptr;
 
diff --git a/cpp/src/pdlp/utilities/ping_pong_graph.cu b/cpp/src/pdlp/utilities/ping_pong_graph.cu
index eb1f31116e..f26dff73eb 100644
--- a/cpp/src/pdlp/utilities/ping_pong_graph.cu
+++ b/cpp/src/pdlp/utilities/ping_pong_graph.cu
@@ -5,13 +5,13 @@
  */
 /* clang-format on */
 
+#include <cuda/stream>
 #include <pdlp/utilities/ping_pong_graph.cuh>
 
 namespace cuopt::mathematical_optimization::pdlp {
 
 template <typename i_t>
-ping_pong_graph_t<i_t>::ping_pong_graph_t(rmm::cuda_stream_view stream_view,
-                                          bool is_legacy_batch_mode)
+ping_pong_graph_t<i_t>::ping_pong_graph_t(cuda::stream_ref stream_view, bool is_legacy_batch_mode)
   : stream_view_(stream_view), is_legacy_batch_mode_(is_legacy_batch_mode)
 {
 }
diff --git a/cpp/src/pdlp/utilities/ping_pong_graph.cuh b/cpp/src/pdlp/utilities/ping_pong_graph.cuh
index 4f895fffd1..683f206880 100644
--- a/cpp/src/pdlp/utilities/ping_pong_graph.cuh
+++ b/cpp/src/pdlp/utilities/ping_pong_graph.cuh
@@ -7,11 +7,10 @@
 
 #pragma once
 
+#include <cuda/stream>
 #include <pdlp/pdlp_constants.hpp>
 #include <utilities/manual_cuda_graph.cuh>
 
-#include <rmm/cuda_stream_view.hpp>
-
 #include <utility>
 
 namespace cuopt::mathematical_optimization::pdlp {
@@ -26,7 +25,7 @@ namespace cuopt::mathematical_optimization::pdlp {
 template <typename i_t>
 class ping_pong_graph_t {
  public:
-  ping_pong_graph_t(rmm::cuda_stream_view stream_view, bool is_legacy_batch_mode = false);
+  ping_pong_graph_t(cuda::stream_ref stream_view, bool is_legacy_batch_mode = false);
   ~ping_pong_graph_t() = default;
 
   // Non-copyable because the underlying manual_cuda_graph_t owns a
@@ -64,7 +63,7 @@ class ping_pong_graph_t {
  private:
   manual_cuda_graph_t even_graph_;
   manual_cuda_graph_t odd_graph_;
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
   bool is_legacy_batch_mode_{false};
 };
 
diff --git a/cpp/src/pdlp/utils.cuh b/cpp/src/pdlp/utils.cuh
index c0a39faa73..ebbc278079 100644
--- a/cpp/src/pdlp/utils.cuh
+++ b/cpp/src/pdlp/utils.cuh
@@ -15,13 +15,13 @@
 #include <random>
 #include <vector>
 
+#include <cuda/stream>
 #include <raft/core/device_span.hpp>
 #include <raft/linalg/binary_op.cuh>
 #include <raft/linalg/detail/cublas_wrappers.hpp>
 #include <raft/linalg/norm.cuh>
 #include <raft/util/cuda_utils.cuh>
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -338,7 +338,7 @@ template <typename f_t>
 void inline compute_sum_bounds_squared(const rmm::device_uvector<f_t>& constraint_lower_bounds,
                                        const rmm::device_uvector<f_t>& constraint_upper_bounds,
                                        rmm::device_scalar<f_t>& out,
-                                       rmm::cuda_stream_view stream_view,
+                                       cuda::stream_ref stream_view,
                                        std::size_t n)
 {
   cuopt_assert(constraint_lower_bounds.size() == constraint_upper_bounds.size(),
@@ -379,7 +379,7 @@ template <typename f_t>
 void inline compute_sum_weighted_squares(const rmm::device_uvector<f_t>& values,
                                          f_t weight,
                                          rmm::device_scalar<f_t>& out,
-                                         rmm::cuda_stream_view stream_view,
+                                         cuda::stream_ref stream_view,
                                          std::size_t n)
 {
   cuopt_assert(n <= values.size(), "n exceeds values size");
@@ -415,7 +415,7 @@ template <typename f_t>
 void inline compute_sum_bounds(const rmm::device_uvector<f_t>& constraint_lower_bounds,
                                const rmm::device_uvector<f_t>& constraint_upper_bounds,
                                f_t* out,
-                               rmm::cuda_stream_view stream_view)
+                               cuda::stream_ref stream_view)
 {
   rmm::device_buffer d_temp_storage;
   size_t bytes = 0;
@@ -449,7 +449,7 @@ template <typename f_t>
 void inline compute_sum_bounds(const rmm::device_uvector<f_t>& constraint_lower_bounds,
                                const rmm::device_uvector<f_t>& constraint_upper_bounds,
                                rmm::device_scalar<f_t>& out,
-                               rmm::cuda_stream_view stream_view)
+                               cuda::stream_ref stream_view)
 {
   compute_sum_bounds(constraint_lower_bounds, constraint_upper_bounds, out.data(), stream_view);
 }
@@ -712,7 +712,7 @@ void inline my_l2_weighted_norm(const f_t* input_vector,
                                 size_t size,
                                 f_t weight,
                                 rmm::device_scalar<f_t>& result,
-                                rmm::cuda_stream_view stream)
+                                cuda::stream_ref stream)
 {
   auto fin_op  = [] __device__(f_t in) { return raft::sqrt(in); };
   auto main_op = [weight] __device__(f_t in, i_t _) { return in * in * weight; };
@@ -732,7 +732,7 @@ template <typename i_t, typename f_t>
 void inline my_l2_weighted_norm(rmm::device_uvector<f_t>& input_vector,
                                 f_t weight,
                                 rmm::device_scalar<f_t>& result,
-                                rmm::cuda_stream_view stream)
+                                cuda::stream_ref stream)
 {
   my_l2_weighted_norm<i_t, f_t>(input_vector.data(), input_vector.size(), weight, result, stream);
 }
diff --git a/cpp/src/routing/assignment.cu b/cpp/src/routing/assignment.cu
index 14862914ef..b63027e673 100644
--- a/cpp/src/routing/assignment.cu
+++ b/cpp/src/routing/assignment.cu
@@ -5,6 +5,7 @@
  */
 /* clang-format on */
 
+#include <cuda/stream>
 #include <cuopt/export.hpp>
 #include <cuopt/routing/assignment.hpp>
 #include <raft/util/cudart_utils.hpp>
@@ -21,7 +22,7 @@ const std::string solution_string_t::empty = "cuOpt solver did not run.";
 const std::string solution_string_t::error = "An error occured while running the cuOpt solver.";
 
 template <typename i_t>
-assignment_t<i_t>::assignment_t(solution_status_t status, rmm::cuda_stream_view stream_view)
+assignment_t<i_t>::assignment_t(solution_status_t status, cuda::stream_ref stream_view)
   : status_(status),
     route_(0, stream_view),
     arrival_stamp_(0, stream_view),
@@ -36,7 +37,7 @@ assignment_t<i_t>::assignment_t(solution_status_t status, rmm::cuda_stream_view
 }
 
 template <typename i_t>
-assignment_t<i_t>::assignment_t(cuopt::logic_error error_status, rmm::cuda_stream_view stream_view)
+assignment_t<i_t>::assignment_t(cuopt::logic_error error_status, cuda::stream_ref stream_view)
   : status_(solution_status_t::ERROR),
     route_(0, stream_view),
     arrival_stamp_(0, stream_view),
@@ -188,7 +189,7 @@ const rmm::device_uvector<i_t>& assignment_t<i_t>::get_accepted() const noexcept
 }
 
 template <typename i_t>
-void assignment_t<i_t>::to_csv(std::string_view filename, rmm::cuda_stream_view stream_view)
+void assignment_t<i_t>::to_csv(std::string_view filename, cuda::stream_ref stream_view)
 {
   std::vector<i_t> route;
   std::vector<double> arrival_stamp;
diff --git a/cpp/src/routing/cpu_routing_problem.cu b/cpp/src/routing/cpu_routing_problem.cu
index f48b10a592..9c8ff346ab 100644
--- a/cpp/src/routing/cpu_routing_problem.cu
+++ b/cpp/src/routing/cpu_routing_problem.cu
@@ -9,6 +9,7 @@
 
 #include <utilities/copy_helpers.hpp>
 
+#include <cuda/stream>
 #include <rmm/device_uvector.hpp>
 
 #include <stdexcept>
@@ -72,14 +73,14 @@ namespace {
 
 template <typename T>
 std::unique_ptr<rmm::device_uvector<T>> copy_vector(std::vector<T> const& host,
-                                                    rmm::cuda_stream_view stream)
+                                                    cuda::stream_ref stream)
 {
   if (host.empty()) { return nullptr; }
   return std::make_unique<rmm::device_uvector<T>>(cuopt::device_copy(host, stream));
 }
 
 std::unique_ptr<rmm::device_uvector<bool>> copy_u8_as_bool(std::vector<uint8_t> const& host,
-                                                           rmm::cuda_stream_view stream)
+                                                           cuda::stream_ref stream)
 {
   if (host.empty()) { return nullptr; }
   std::vector<bool> as_bool(host.begin(), host.end());
diff --git a/cpp/src/routing/crossovers/ox_graph.hpp b/cpp/src/routing/crossovers/ox_graph.hpp
index 55d2e39a0b..4783c542b1 100644
--- a/cpp/src/routing/crossovers/ox_graph.hpp
+++ b/cpp/src/routing/crossovers/ox_graph.hpp
@@ -7,13 +7,15 @@
 
 #pragma once
 
+#include <cuda/stream>
+
 namespace cuopt {
 namespace routing {
 namespace detail {
 
 template <typename i_t, typename f_t>
 struct ox_graph_t {
-  ox_graph_t(i_t n_buckets_, i_t size, i_t max_nodes_per_row, rmm::cuda_stream_view stream)
+  ox_graph_t(i_t n_buckets_, i_t size, i_t max_nodes_per_row, cuda::stream_ref stream)
     : row_sizes(n_buckets_ * size, stream),
       route_ids(n_buckets_ * size, stream),
       // allocate with the max size
@@ -39,7 +41,7 @@ struct ox_graph_t {
     std::vector<int> buckets;
   };
 
-  host_t to_host(rmm::cuda_stream_view stream)
+  host_t to_host(cuda::stream_ref stream)
   {
     host_t h;
     h.row_sizes = host_copy(row_sizes, stream);
@@ -50,7 +52,7 @@ struct ox_graph_t {
     return h;
   }
 
-  void resize(i_t n_buckets_, i_t size, i_t max_nodes_per_row, rmm::cuda_stream_view stream)
+  void resize(i_t n_buckets_, i_t size, i_t max_nodes_per_row, cuda::stream_ref stream)
   {
     n_buckets = n_buckets_;
     row_sizes.resize(n_buckets * size, stream);
diff --git a/cpp/src/routing/crossovers/ox_recombiner.cuh b/cpp/src/routing/crossovers/ox_recombiner.cuh
index f16f9d2a11..e7b6cf04ec 100644
--- a/cpp/src/routing/crossovers/ox_recombiner.cuh
+++ b/cpp/src/routing/crossovers/ox_recombiner.cuh
@@ -17,6 +17,7 @@
 #include <random>
 #include <vector>
 
+#include <cuda/stream>
 #include <rmm/device_uvector.hpp>
 
 #include <cub/cub.cuh>
@@ -94,7 +95,7 @@ struct OX {
   ox_graph_t<int, float> d_graph;
   ox_graph_t<int, float> transpose_graph;
 
-  explicit OX(size_t nodes_number, const costs& weight, rmm::cuda_stream_view stream_view)
+  explicit OX(size_t nodes_number, const costs& weight, cuda::stream_ref stream_view)
     : mt(rd()),
       problem_size(nodes_number),
       graph(problem_size),
@@ -518,7 +519,7 @@ struct OX {
     }
   }
 
-  void test_transpose_graph(rmm::cuda_stream_view stream)
+  void test_transpose_graph(cuda::stream_ref stream)
   {
     std::vector<std::vector<std::tuple<int, double, int>>> h_transpose_graph(offspring.size());
     for (size_t i = 0; i < h_transpose_graph.size(); ++i) {
@@ -846,7 +847,7 @@ struct OX {
   }
 
   void adj_to_host(std::vector<std::vector<std::tuple<int, double, int>>>& h_graph,
-                   rmm::cuda_stream_view stream)
+                   cuda::stream_ref stream)
   {
     auto tmp_graph = d_graph.to_host(stream);
     for (int veh = 0; veh < n_buckets; ++veh) {
diff --git a/cpp/src/routing/cuda_graph.cuh b/cpp/src/routing/cuda_graph.cuh
index 40e1b95f49..b260697e02 100644
--- a/cpp/src/routing/cuda_graph.cuh
+++ b/cpp/src/routing/cuda_graph.cuh
@@ -5,11 +5,10 @@
  */
 /* clang-format on */
 
+#include <cuda/stream>
 #include <cuopt/error.hpp>
 #include <utilities/macros.cuh>
 
-#include <rmm/cuda_stream_view.hpp>
-
 #pragma once
 
 namespace cuopt {
@@ -18,7 +17,7 @@ namespace detail {
 
 // This is not a thread-safe class, be careful on multi-threading
 struct cuda_graph_t {
-  void start_capture(rmm::cuda_stream_view stream)
+  void start_capture(cuda::stream_ref stream)
   {
     // Use ThreadLocal mode to allow multi-threaded batch execution
     // Global mode blocks other streams from performing operations during capture
@@ -26,7 +25,7 @@ struct cuda_graph_t {
     capture_started = true;
   }
 
-  void end_capture(rmm::cuda_stream_view stream)
+  void end_capture(cuda::stream_ref stream)
   {
     cuopt_assert(capture_started, "start_capture was not called before end_capture!");
     cuopt_expects(capture_started, error_type_t::RuntimeError, "A runtime error occurred!");
@@ -52,7 +51,7 @@ struct cuda_graph_t {
     cudaGraphDestroy(graph);
   }
 
-  void launch_graph(rmm::cuda_stream_view stream) { cudaGraphLaunch(instance, stream.get()); }
+  void launch_graph(cuda::stream_ref stream) { cudaGraphLaunch(instance, stream.get()); }
 
   bool graph_created   = false;
   bool capture_started = false;
diff --git a/cpp/src/routing/fleet_info.hpp b/cpp/src/routing/fleet_info.hpp
index 1a37c66554..a40fefc04e 100644
--- a/cpp/src/routing/fleet_info.hpp
+++ b/cpp/src/routing/fleet_info.hpp
@@ -7,6 +7,7 @@
 
 #pragma once
 
+#include <cuda/stream>
 #include <routing/routing_details.hpp>
 #include <routing/utilities/md_utils.hpp>
 #include "fleet_order_constraints.hpp"
@@ -55,7 +56,7 @@ class fleet_info_t {
 
   constexpr bool is_homogenous() const { return is_homogenous_; }
 
-  void resize(i_t size, rmm::cuda_stream_view stream)
+  void resize(i_t size, cuda::stream_ref stream)
   {
     v_earliest_time_.resize(size, stream);
     v_latest_time_.resize(size, stream);
@@ -70,7 +71,7 @@ class fleet_info_t {
     v_buckets_.resize(size, stream);
   }
 
-  auto to_host(rmm::cuda_stream_view stream)
+  auto to_host(cuda::stream_ref stream)
   {
     host_t h;
     h.break_offset            = host_copy(v_break_offset_, stream);
@@ -237,7 +238,7 @@ class fleet_info_t {
 
   constexpr raft::device_span<i_t const> get_break_vector(i_t truck_id,
                                                           const rmm::device_uvector<i_t>& vec,
-                                                          rmm::cuda_stream_view stream) const
+                                                          cuda::stream_ref stream) const
   {
     if (!vec.is_empty()) {
       i_t offset    = v_break_offset_.element(truck_id, stream);
@@ -248,8 +249,7 @@ class fleet_info_t {
     }
   }
 
-  constexpr VehicleInfo<f_t> get_vehicle_info(const i_t vehicle_id,
-                                              rmm::cuda_stream_view stream) const
+  constexpr VehicleInfo<f_t> get_vehicle_info(const i_t vehicle_id, cuda::stream_ref stream) const
   {
     return v_vehicle_infos_.element(vehicle_id, stream);
   }
diff --git a/cpp/src/routing/fleet_order_constraints.hpp b/cpp/src/routing/fleet_order_constraints.hpp
index c6be63a87c..e28073e171 100644
--- a/cpp/src/routing/fleet_order_constraints.hpp
+++ b/cpp/src/routing/fleet_order_constraints.hpp
@@ -10,6 +10,7 @@
 #include <cuopt/routing/data_model_view.hpp>
 #include <utilities/copy_helpers.hpp>
 
+#include <cuda/stream>
 #include <raft/core/span.hpp>
 
 #include <thrust/fill.h>
@@ -67,7 +68,7 @@ struct fleet_order_constraints_t {
     i_t n_vehicles;
   };
 
-  host_t to_host(rmm::cuda_stream_view stream)
+  host_t to_host(cuda::stream_ref stream)
   {
     host_t h;
     h.order_service_times = host_copy(order_service_times, stream);
diff --git a/cpp/src/routing/ges/compute_fragment_ejections.cu b/cpp/src/routing/ges/compute_fragment_ejections.cu
index 46db0c0cbb..c96b06d6e1 100644
--- a/cpp/src/routing/ges/compute_fragment_ejections.cu
+++ b/cpp/src/routing/ges/compute_fragment_ejections.cu
@@ -5,6 +5,7 @@
  */
 /* clang-format on */
 
+#include <cuda/stream>
 #include <utilities/cuda_helpers.cuh>
 #include "../solution/solution.cuh"
 #include "compute_delivery_insertions.cuh"
@@ -122,7 +123,7 @@ bool set_shmem_for_kernel_get_best_insertion_ejection_solution(size_t dynamic_sh
 
 template <int BLOCK_SIZE, typename i_t, typename f_t, request_t REQUEST>
 void launch_kernel_get_best_insertion_ejection_solution(
-  dim3 grid, dim3 blocks, size_t shmem_bytes, void** kernel_args, rmm::cuda_stream_view stream)
+  dim3 grid, dim3 blocks, size_t shmem_bytes, void** kernel_args, cuda::stream_ref stream)
 {
   RAFT_CUDA_TRY(cudaLaunchKernel(
     (void*)kernel_get_best_insertion_ejection_solution<BLOCK_SIZE, i_t, f_t, REQUEST>,
@@ -141,7 +142,7 @@ void launch_kernel_get_best_insertion_ejection_solution(
     size_t dynamic_shmem_size);                                                               \
   template void                                                                               \
   launch_kernel_get_best_insertion_ejection_solution<BLOCK_SIZE, int, float, request_t::REQ>( \
-    dim3 grid, dim3 blocks, size_t shmem_bytes, void** kernel_args, rmm::cuda_stream_view stream);
+    dim3 grid, dim3 blocks, size_t shmem_bytes, void** kernel_args, cuda::stream_ref stream);
 
 CUOPT_INSTANTIATE_GET_BEST_INSERTION_EJECTION(32, PDP)
 CUOPT_INSTANTIATE_GET_BEST_INSERTION_EJECTION(64, PDP)
diff --git a/cpp/src/routing/ges/compute_fragment_ejections.cuh b/cpp/src/routing/ges/compute_fragment_ejections.cuh
index 6e560474d1..2dd599e951 100644
--- a/cpp/src/routing/ges/compute_fragment_ejections.cuh
+++ b/cpp/src/routing/ges/compute_fragment_ejections.cuh
@@ -7,6 +7,7 @@
 
 #pragma once
 
+#include <cuda/stream>
 #include <utilities/cuda_helpers.cuh>
 #include "../solution/solution.cuh"
 #include "found_solution.cuh"
@@ -40,7 +41,7 @@ bool set_shmem_for_kernel_get_best_insertion_ejection_solution(size_t dynamic_sh
 
 template <int BLOCK_SIZE, typename i_t, typename f_t, request_t REQUEST>
 void launch_kernel_get_best_insertion_ejection_solution(
-  dim3 grid, dim3 blocks, size_t shmem_bytes, void** kernel_args, rmm::cuda_stream_view stream);
+  dim3 grid, dim3 blocks, size_t shmem_bytes, void** kernel_args, cuda::stream_ref stream);
 
 template <int BLOCK_SIZE,
           typename i_t,
diff --git a/cpp/src/routing/ges/ejection_pool.cuh b/cpp/src/routing/ges/ejection_pool.cuh
index b07160accd..9cf975c28f 100644
--- a/cpp/src/routing/ges/ejection_pool.cuh
+++ b/cpp/src/routing/ges/ejection_pool.cuh
@@ -13,7 +13,7 @@
 #include <routing/utilities/seed_generator.cuh>
 #include <utilities/cuda_helpers.cuh>
 
-#include <rmm/cuda_stream_view.hpp>
+#include <cuda/stream>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -41,7 +41,7 @@ __global__ static void device_random_shuffle(elemt_t* data, int size, int64_t se
  */
 template <class elemt_t, int max_stack_size = std::numeric_limits<int>::max()>
 struct ejection_pool_t {
-  ejection_pool_t(int max_ejection_pool_size, rmm::cuda_stream_view stream)
+  ejection_pool_t(int max_ejection_pool_size, cuda::stream_ref stream)
     : stack_(max_ejection_pool_size, stream), index_(-1), stream_(stream)
   {
   }
@@ -154,7 +154,7 @@ struct ejection_pool_t {
 
   rmm::device_uvector<elemt_t> stack_;
   int index_;
-  rmm::cuda_stream_view stream_;
+  cuda::stream_ref stream_;
   std::uniform_int_distribution<int> dist{0, std::numeric_limits<int>::max()};
   std::mt19937 gen{66742};
 };
diff --git a/cpp/src/routing/local_search/cycle_finder/cycle.hpp b/cpp/src/routing/local_search/cycle_finder/cycle.hpp
index e6a8aec57b..e0d19214e2 100644
--- a/cpp/src/routing/local_search/cycle_finder/cycle.hpp
+++ b/cpp/src/routing/local_search/cycle_finder/cycle.hpp
@@ -12,8 +12,8 @@
 #include <utilities/vector_helpers.cuh>
 #include "../../solution/solution_handle.cuh"
 
+#include <cuda/stream>
 #include <raft/core/device_span.hpp>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -23,7 +23,7 @@ namespace detail {
 
 template <typename i_t, typename f_t>
 struct ret_cycles_t {
-  ret_cycles_t(size_t max_size, rmm::cuda_stream_view stream_view)
+  ret_cycles_t(size_t max_size, cuda::stream_ref stream_view)
     : paths(max_size, stream_view),
       offsets(max_size, stream_view),
       n_cycles_(zero_v<i_t>, stream_view),
@@ -46,7 +46,7 @@ struct ret_cycles_t {
     i_t n_cycles;
   };
 
-  host_t to_host(rmm::cuda_stream_view stream)
+  host_t to_host(cuda::stream_ref stream)
   {
     host_t h;
     h.paths    = host_copy(paths, stream);
diff --git a/cpp/src/routing/local_search/cycle_finder/cycle_finder.cu b/cpp/src/routing/local_search/cycle_finder/cycle_finder.cu
index c9994f672c..d8d2d0f3c0 100644
--- a/cpp/src/routing/local_search/cycle_finder/cycle_finder.cu
+++ b/cpp/src/routing/local_search/cycle_finder/cycle_finder.cu
@@ -7,6 +7,7 @@
 
 #include "cycle_finder_kernels.cuh"
 
+#include <cuda/stream>
 #include <raft/core/nvtx.hpp>
 
 #include <thrust/iterator/constant_iterator.h>
@@ -138,7 +139,7 @@ bool ExactCycleFinder<i_t, f_t, max_routes>::call_find(graph_t<i_t, f_t>& graph,
 }
 
 template <typename map_key_t, typename value_t>
-void detail::device_map_t<map_key_t, value_t>::clear(rmm::cuda_stream_view stream)
+void detail::device_map_t<map_key_t, value_t>::clear(cuda::stream_ref stream)
 {
   auto max_vals  = max_level * max_available;
   auto n_threads = 256;
@@ -149,7 +150,7 @@ void detail::device_map_t<map_key_t, value_t>::clear(rmm::cuda_stream_view strea
 
 template <size_t max_routes>
 bool test_empty(typename detail::device_map_t<key_t<max_routes>, double>::view_t const map_view,
-                rmm::cuda_stream_view stream)
+                cuda::stream_ref stream)
 {
   auto max_vals  = map_view.max_available;
   auto n_threads = 256;
diff --git a/cpp/src/routing/local_search/cycle_finder/cycle_finder.hpp b/cpp/src/routing/local_search/cycle_finder/cycle_finder.hpp
index 7db5d0c417..d94cddf39e 100644
--- a/cpp/src/routing/local_search/cycle_finder/cycle_finder.hpp
+++ b/cpp/src/routing/local_search/cycle_finder/cycle_finder.hpp
@@ -14,8 +14,8 @@
 
 #include "../../solution/solution_handle.cuh"
 
+#include <cuda/stream>
 #include <raft/util/cudart_utils.hpp>
-#include <rmm/cuda_stream_view.hpp>
 
 namespace cuopt {
 namespace routing {
@@ -60,7 +60,7 @@ struct path_t {
     reset(handle_ptr_->get_stream());
   }
 
-  void reset(rmm::cuda_stream_view stream)
+  void reset(cuda::stream_ref stream)
   {
     n_cycles.set_value_to_zero_async(stream);
     all_found.set_value_to_zero_async(stream);
@@ -112,7 +112,7 @@ struct path_t {
 
 template <typename i_t, typename f_t, size_t max_routes>
 struct cycle_candidates_t {
-  cycle_candidates_t(size_t size_, int n_paths_, rmm::cuda_stream_view stream)
+  cycle_candidates_t(size_t size_, int n_paths_, cuda::stream_ref stream)
     : keys(size_ * n_paths_, stream),
       costs(size_ * n_paths_, stream),
       level_vec(size_ * n_paths_, stream),
diff --git a/cpp/src/routing/local_search/cycle_finder/cycle_graph.hpp b/cpp/src/routing/local_search/cycle_finder/cycle_graph.hpp
index 3c28f78bc4..3bdbef0567 100644
--- a/cpp/src/routing/local_search/cycle_finder/cycle_graph.hpp
+++ b/cpp/src/routing/local_search/cycle_finder/cycle_graph.hpp
@@ -13,6 +13,7 @@
 #include <utilities/copy_helpers.hpp>
 #include <utilities/vector_helpers.cuh>
 
+#include <cuda/stream>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -29,7 +30,7 @@ constexpr int max_graph_nodes_per_row = 1024;
 
 template <typename i_t, typename f_t>
 struct graph_t {
-  graph_t(i_t size, rmm::cuda_stream_view stream)
+  graph_t(i_t size, cuda::stream_ref stream)
     : row_sizes(size, stream),
       route_ids(size, stream),
       // allocate with the max size
@@ -45,7 +46,7 @@ struct graph_t {
     std::vector<double> weights;
   };
 
-  host_t to_host(rmm::cuda_stream_view stream)
+  host_t to_host(cuda::stream_ref stream)
   {
     host_t h;
     h.row_sizes = host_copy(row_sizes, stream);
diff --git a/cpp/src/routing/local_search/cycle_finder/device_map.cuh b/cpp/src/routing/local_search/cycle_finder/device_map.cuh
index 3a2b7a31f4..d6745c4cb6 100644
--- a/cpp/src/routing/local_search/cycle_finder/device_map.cuh
+++ b/cpp/src/routing/local_search/cycle_finder/device_map.cuh
@@ -1,6 +1,6 @@
 /* clang-format off */
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 /* clang-format on */
@@ -13,6 +13,7 @@
 
 #include <utilities/macros.cuh>
 
+#include <cuda/stream>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -58,7 +59,7 @@ struct device_map_t {
   {
   }
 
-  void clear(rmm::cuda_stream_view stream);
+  void clear(cuda::stream_ref stream);
 
   uint32_t get_max_size() const
   {
@@ -66,7 +67,7 @@ struct device_map_t {
     return adj_max_size;
   }
 
-  size_t get_size(int level, rmm::cuda_stream_view stream) const
+  size_t get_size(int level, cuda::stream_ref stream) const
   {
     return std::min(get_max_size(), occupied.element(level, stream));
   }
diff --git a/cpp/src/routing/local_search/hvrp/vehicle_assignment.cuh b/cpp/src/routing/local_search/hvrp/vehicle_assignment.cuh
index da8fef1258..42369093c0 100644
--- a/cpp/src/routing/local_search/hvrp/vehicle_assignment.cuh
+++ b/cpp/src/routing/local_search/hvrp/vehicle_assignment.cuh
@@ -1,12 +1,13 @@
 /* clang-format off */
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 /* clang-format on */
 
 #pragma once
 
+#include <cuda/stream>
 #include "../../solution/solution.cuh"
 #include "../move_candidates/move_candidates.cuh"
 
@@ -59,7 +60,7 @@ struct vehicle_assignment_t {
     gl_lock.set_value_to_zero_async(sol_handle_->get_stream());
   }
 
-  void resize(i_t n_routes, i_t n_buckets, rmm::cuda_stream_view stream_view)
+  void resize(i_t n_routes, i_t n_buckets, cuda::stream_ref stream_view)
   {
     k_regrets   = std::min(n_buckets, k_max_regrets);
     auto k_iter = k_regrets - 1;
diff --git a/cpp/src/routing/order_info.hpp b/cpp/src/routing/order_info.hpp
index d20c46a8ed..b7e8ba1e11 100644
--- a/cpp/src/routing/order_info.hpp
+++ b/cpp/src/routing/order_info.hpp
@@ -7,6 +7,7 @@
 
 #pragma once
 
+#include <cuda/stream>
 #include <raft/core/device_span.hpp>
 #include <routing/routing_details.hpp>
 #include <routing/structures.hpp>
@@ -43,7 +44,7 @@ class order_info_t {
     return is_pdp() ? get_num_depot_excluded_orders() / 2 : get_num_depot_excluded_orders();
   }
 
-  void resize(i_t size, bool is_pickup, rmm::cuda_stream_view stream)
+  void resize(i_t size, bool is_pickup, cuda::stream_ref stream)
   {
     v_demand_.resize(size, stream);
     v_earliest_time_.resize(size, stream);
@@ -57,7 +58,7 @@ class order_info_t {
 
   bool is_pdp() const { return !v_pair_indices_.is_empty(); }
 
-  auto to_host(rmm::cuda_stream_view stream)
+  auto to_host(cuda::stream_ref stream)
   {
     host_t h;
     h.earliest_time   = cuopt::host_copy(v_earliest_time_, stream);
diff --git a/cpp/src/routing/route/break_route.cuh b/cpp/src/routing/route/break_route.cuh
index 1d5b3472f9..f6acaa29bf 100644
--- a/cpp/src/routing/route/break_route.cuh
+++ b/cpp/src/routing/route/break_route.cuh
@@ -12,6 +12,7 @@
 #include "../solution/solution_handle.cuh"
 #include "routing/routing_helpers.cuh"
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 #include <raft/core/nvtx.hpp>
 
@@ -44,7 +45,7 @@ class break_route_t {
 
   break_route_t& operator=(break_route_t&& break_route) = default;
 
-  void resize(i_t max_nodes_per_route, rmm::cuda_stream_view stream)
+  void resize(i_t max_nodes_per_route, cuda::stream_ref stream)
   {
     breaks_forward.resize(max_nodes_per_route, stream);
     breaks_backward.resize(max_nodes_per_route, stream);
diff --git a/cpp/src/routing/route/capacity_route.cuh b/cpp/src/routing/route/capacity_route.cuh
index 776262a497..b91acbd30a 100644
--- a/cpp/src/routing/route/capacity_route.cuh
+++ b/cpp/src/routing/route/capacity_route.cuh
@@ -13,6 +13,7 @@
 #include "../solution/solution_handle.cuh"
 #include "routing/routing_helpers.cuh"
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 #include <raft/core/nvtx.hpp>
 
@@ -53,7 +54,7 @@ class capacity_route_t {
 
   capacity_route_t& operator=(capacity_route_t&& capacity_route) = default;
 
-  void resize(i_t max_nodes_per_route, rmm::cuda_stream_view stream)
+  void resize(i_t max_nodes_per_route, cuda::stream_ref stream)
   {
     i_t n_dims = dim_info.n_capacity_dimensions;
     if (n_dims == 0) { return; }
diff --git a/cpp/src/routing/route/distance_route.cuh b/cpp/src/routing/route/distance_route.cuh
index a5f98c13ce..158283a850 100644
--- a/cpp/src/routing/route/distance_route.cuh
+++ b/cpp/src/routing/route/distance_route.cuh
@@ -12,6 +12,7 @@
 #include "../solution/solution_handle.cuh"
 #include "routing/routing_helpers.cuh"
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 #include <raft/core/nvtx.hpp>
 
@@ -47,7 +48,7 @@ class distance_route_t {
 
   distance_route_t& operator=(distance_route_t&& distance_route) = default;
 
-  void resize(i_t max_nodes_per_route, rmm::cuda_stream_view stream)
+  void resize(i_t max_nodes_per_route, cuda::stream_ref stream)
   {
     distance_forward.resize(max_nodes_per_route, stream);
     distance_backward.resize(max_nodes_per_route, stream);
diff --git a/cpp/src/routing/route/mismatch_route.cuh b/cpp/src/routing/route/mismatch_route.cuh
index 78975750e0..10a0a801b7 100644
--- a/cpp/src/routing/route/mismatch_route.cuh
+++ b/cpp/src/routing/route/mismatch_route.cuh
@@ -11,6 +11,7 @@
 #include "../node/mismatch_node.cuh"
 #include "../solution/solution_handle.cuh"
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
 #include <rmm/device_uvector.hpp>
@@ -42,7 +43,7 @@ class mismatch_route_t {
 
   mismatch_route_t& operator=(mismatch_route_t&& mismatch_route) = default;
 
-  void resize(i_t max_nodes_per_route, rmm::cuda_stream_view stream)
+  void resize(i_t max_nodes_per_route, cuda::stream_ref stream)
   {
     mismatch_forward.resize(max_nodes_per_route, stream);
     mismatch_backward.resize(max_nodes_per_route, stream);
diff --git a/cpp/src/routing/route/pdp_route.cuh b/cpp/src/routing/route/pdp_route.cuh
index dd20e2fec3..449a2e182f 100644
--- a/cpp/src/routing/route/pdp_route.cuh
+++ b/cpp/src/routing/route/pdp_route.cuh
@@ -12,6 +12,7 @@
 #include "../solution/solution_handle.cuh"
 #include "tsp_route.cuh"
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 #include <raft/core/nvtx.hpp>
 
@@ -60,7 +61,7 @@ class request_route_t<i_t, f_t, REQUEST, std::enable_if_t<REQUEST == request_t::
     std::cout << "]\n";
   }
 
-  void resize(i_t max_nodes_per_route, bool is_tsp, rmm::cuda_stream_view stream)
+  void resize(i_t max_nodes_per_route, bool is_tsp, cuda::stream_ref stream)
   {
     node_info.resize(max_nodes_per_route, stream);
     brother_info.resize(max_nodes_per_route, stream);
@@ -182,7 +183,7 @@ class request_route_t<i_t, f_t, REQUEST, std::enable_if_t<REQUEST == request_t::
     std::cout << "]\n";
   }
 
-  void resize(i_t max_nodes_per_route, bool is_tsp, rmm::cuda_stream_view stream)
+  void resize(i_t max_nodes_per_route, bool is_tsp, cuda::stream_ref stream)
   {
     node_info.resize(max_nodes_per_route, stream);
     if (is_tsp) { tsp_requests.resize(max_nodes_per_route, stream); }
diff --git a/cpp/src/routing/route/prize_route.cuh b/cpp/src/routing/route/prize_route.cuh
index 80b27061b5..dd4ab62c7d 100644
--- a/cpp/src/routing/route/prize_route.cuh
+++ b/cpp/src/routing/route/prize_route.cuh
@@ -12,6 +12,7 @@
 #include "../solution/solution_handle.cuh"
 #include "routing/routing_helpers.cuh"
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 #include <raft/core/nvtx.hpp>
 
@@ -46,7 +47,7 @@ class prize_route_t {
 
   prize_route_t& operator=(prize_route_t&& prize_route) = default;
 
-  void resize(i_t max_nodes_per_route, rmm::cuda_stream_view stream)
+  void resize(i_t max_nodes_per_route, cuda::stream_ref stream)
   {
     prize.resize(max_nodes_per_route, stream);
     prize_forward.resize(max_nodes_per_route, stream);
diff --git a/cpp/src/routing/route/service_time_route.cuh b/cpp/src/routing/route/service_time_route.cuh
index 03c48b2e42..46f433db4e 100644
--- a/cpp/src/routing/route/service_time_route.cuh
+++ b/cpp/src/routing/route/service_time_route.cuh
@@ -12,6 +12,7 @@
 #include "../solution/solution_handle.cuh"
 #include "routing/routing_helpers.cuh"
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -42,7 +43,7 @@ class service_time_route_t {
 
   service_time_route_t& operator=(service_time_route_t&& service_time_route) = default;
 
-  void resize(i_t max_nodes_per_route, rmm::cuda_stream_view stream)
+  void resize(i_t max_nodes_per_route, cuda::stream_ref stream)
   {
     service_time_forward.resize(max_nodes_per_route, stream);
     service_time_backward.resize(max_nodes_per_route, stream);
diff --git a/cpp/src/routing/route/tasks_route.cuh b/cpp/src/routing/route/tasks_route.cuh
index 3624d647e7..fd6d40b099 100644
--- a/cpp/src/routing/route/tasks_route.cuh
+++ b/cpp/src/routing/route/tasks_route.cuh
@@ -11,6 +11,7 @@
 #include "../node/tasks_node.cuh"
 #include "../solution/solution_handle.cuh"
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
 #include <rmm/device_uvector.hpp>
@@ -40,7 +41,7 @@ class tasks_route_t {
 
   tasks_route_t& operator=(tasks_route_t&& tasks_route) = default;
 
-  void resize(i_t max_nodes_per_route, rmm::cuda_stream_view stream)
+  void resize(i_t max_nodes_per_route, cuda::stream_ref stream)
   {
     tasks_forward.resize(max_nodes_per_route, stream);
     tasks_backward.resize(max_nodes_per_route, stream);
diff --git a/cpp/src/routing/route/time_route.cuh b/cpp/src/routing/route/time_route.cuh
index 21448c4273..92d0884e5f 100644
--- a/cpp/src/routing/route/time_route.cuh
+++ b/cpp/src/routing/route/time_route.cuh
@@ -12,6 +12,7 @@
 #include "../solution/solution_handle.cuh"
 #include "routing/routing_helpers.cuh"
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 #include <raft/core/nvtx.hpp>
 
@@ -66,7 +67,7 @@ class time_route_t {
 
   time_route_t& operator=(time_route_t&& time_route) = default;
 
-  void resize(i_t max_nodes_per_route, rmm::cuda_stream_view stream)
+  void resize(i_t max_nodes_per_route, cuda::stream_ref stream)
   {
     departure_forward.resize(max_nodes_per_route, stream);
     excess_forward.resize(max_nodes_per_route, stream);
diff --git a/cpp/src/routing/route/tsp_route.cuh b/cpp/src/routing/route/tsp_route.cuh
index 9b7eeeee56..310418e7ea 100644
--- a/cpp/src/routing/route/tsp_route.cuh
+++ b/cpp/src/routing/route/tsp_route.cuh
@@ -11,6 +11,7 @@
 #include "../node/node.cuh"
 #include "../solution/solution_handle.cuh"
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 #include <raft/core/nvtx.hpp>
 
@@ -40,7 +41,7 @@ class tsp_route_t {
 
   tsp_route_t& operator=(tsp_route_t&& tsp_route) = default;
 
-  void resize(i_t max_nodes_per_route, rmm::cuda_stream_view stream)
+  void resize(i_t max_nodes_per_route, cuda::stream_ref stream)
   {
     pred.resize(max_nodes_per_route, stream);
     succ.resize(max_nodes_per_route, stream);
diff --git a/cpp/src/routing/route/vehicle_fixed_cost_route.cuh b/cpp/src/routing/route/vehicle_fixed_cost_route.cuh
index 1e246fbb6e..45222eab17 100644
--- a/cpp/src/routing/route/vehicle_fixed_cost_route.cuh
+++ b/cpp/src/routing/route/vehicle_fixed_cost_route.cuh
@@ -12,6 +12,7 @@
 #include "../solution/solution_handle.cuh"
 #include "routing/routing_helpers.cuh"
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 
 #include <thrust/tuple.h>
@@ -38,10 +39,7 @@ class vehicle_fixed_cost_route_t {
   vehicle_fixed_cost_route_t& operator=(vehicle_fixed_cost_route_t&& vehicle_fixed_cost_route) =
     default;
 
-  void resize([[maybe_unused]] i_t max_nodes_per_route,
-              [[maybe_unused]] rmm::cuda_stream_view stream)
-  {
-  }
+  void resize([[maybe_unused]] i_t max_nodes_per_route, [[maybe_unused]] cuda::stream_ref stream) {}
 
   struct view_t {
     DI vehicle_fixed_cost_node_t<i_t, f_t> get_node(i_t idx) const
diff --git a/cpp/src/routing/solution/pool_allocator.cuh b/cpp/src/routing/solution/pool_allocator.cuh
index 393740c351..cc7f322a99 100644
--- a/cpp/src/routing/solution/pool_allocator.cuh
+++ b/cpp/src/routing/solution/pool_allocator.cuh
@@ -7,6 +7,7 @@
 
 #pragma once
 
+#include <cuda/stream>
 #include "solution.cuh"
 
 #include "../ges/guided_ejection_search.cuh"
@@ -44,7 +45,7 @@ class pool_allocator_t {
  public:
   pool_allocator_t(const Problem& problem_,
                    i_t n_solutions_,
-                   rmm::cuda_stream_view stream_,
+                   cuda::stream_ref stream_,
                    i_t desired_n_routes = -1)
     : problem(problem_), stream(stream_)
   {
@@ -73,7 +74,7 @@ class pool_allocator_t {
   void sync_all_streams() const { stream.sync(); }
 
   // problem description
-  rmm::cuda_stream_view stream;
+  cuda::stream_ref stream;
   const Problem& problem;
   std::vector<std::unique_ptr<solution_handle_t<i_t, f_t>>> sol_handles;
   // keep a thread safe pool of local search and ges objects that can be reused
diff --git a/cpp/src/routing/solution/route_node_map.cuh b/cpp/src/routing/solution/route_node_map.cuh
index a4a1b171aa..b7af5867e0 100644
--- a/cpp/src/routing/solution/route_node_map.cuh
+++ b/cpp/src/routing/solution/route_node_map.cuh
@@ -9,6 +9,7 @@
 
 #include <thrust/pair.h>
 #include <thrust/tuple.h>
+#include <cuda/stream>
 #include <raft/core/device_span.hpp>
 #include <raft/util/cuda_utils.cuh>
 #include <rmm/device_uvector.hpp>
@@ -20,7 +21,7 @@ namespace detail {
 template <typename i_t>
 class route_node_map_t {
  public:
-  route_node_map_t(const int num_orders, rmm::cuda_stream_view stream)
+  route_node_map_t(const int num_orders, cuda::stream_ref stream)
     : route_id_per_node(num_orders, stream), intra_route_idx_per_node(num_orders, stream)
   {
     thrust::fill(rmm::exec_policy(stream), route_id_per_node.begin(), route_id_per_node.end(), -1);
@@ -30,13 +31,13 @@ class route_node_map_t {
                  -1);
   }
 
-  route_node_map_t(const route_node_map_t& other, rmm::cuda_stream_view stream)
+  route_node_map_t(const route_node_map_t& other, cuda::stream_ref stream)
     : route_id_per_node(other.route_id_per_node, stream),
       intra_route_idx_per_node(other.intra_route_idx_per_node, stream)
   {
   }
 
-  void copy_from(const route_node_map_t& other, rmm::cuda_stream_view stream)
+  void copy_from(const route_node_map_t& other, cuda::stream_ref stream)
   {
     raft::copy(intra_route_idx_per_node.data(),
                other.intra_route_idx_per_node.data(),
diff --git a/cpp/src/routing/solution/solution_handle.cuh b/cpp/src/routing/solution/solution_handle.cuh
index 38675021b5..e2572d5ca3 100644
--- a/cpp/src/routing/solution/solution_handle.cuh
+++ b/cpp/src/routing/solution/solution_handle.cuh
@@ -7,8 +7,8 @@
 
 #pragma once
 
+#include <cuda/stream>
 #include <raft/util/cudart_utils.hpp>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 
 #include <mutex>
@@ -30,7 +30,7 @@ class solution_handle_t {
   solution_handle_t(solution_handle_t&&)                 = delete;
   solution_handle_t& operator=(solution_handle_t&&)      = delete;
 
-  solution_handle_t(rmm::cuda_stream_view stream)
+  solution_handle_t(cuda::stream_ref stream)
     : dev_id_([]() -> i_t {
         i_t cur_dev = -1;
         RAFT_CUDA_TRY(cudaGetDevice(&cur_dev));
@@ -42,7 +42,7 @@ class solution_handle_t {
   }
 
   rmm::exec_policy& get_thrust_policy() const noexcept { return *thrust_policy_; }
-  rmm::cuda_stream_view get_stream() const noexcept { return stream_view_; }
+  cuda::stream_ref get_stream() const noexcept { return stream_view_; }
   i_t get_device() const { return dev_id_; }
   void sync_stream() const { stream_view_.sync(); };
 
@@ -72,7 +72,7 @@ class solution_handle_t {
   mutable bool device_prop_initialized_{false};
 
   mutable bool shared_attr_initialized_{false};
-  rmm::cuda_stream_view stream_view_{};
+  cuda::stream_ref stream_view_{};
   // this is a shared pointer to be able to copy construct and keep a copy of a solution
   std::shared_ptr<rmm::exec_policy> thrust_policy_{nullptr};
 };
diff --git a/cpp/src/routing/util_kernels/runtime_checks.cu b/cpp/src/routing/util_kernels/runtime_checks.cu
index b9f28f6a18..35305b4004 100644
--- a/cpp/src/routing/util_kernels/runtime_checks.cu
+++ b/cpp/src/routing/util_kernels/runtime_checks.cu
@@ -5,6 +5,7 @@
  */
 /* clang-format on */
 
+#include <cuda/stream>
 #include <utilities/vector_helpers.cuh>
 #include "../solution/solution.cuh"
 
@@ -246,7 +247,7 @@ __global__ void check_breaks(typename solution_t<i_t, f_t, REQUEST>::view_t solu
 
 template <typename i_t, typename f_t, request_t REQUEST>
 bool global_runtime_checks_(solution_t<i_t, f_t, REQUEST>& solution,
-                            rmm::cuda_stream_view stream,
+                            cuda::stream_ref stream,
                             bool all_nodes_should_be_served,
                             bool check_feasible)
 {
diff --git a/cpp/src/routing/utilities/check_input.cu b/cpp/src/routing/utilities/check_input.cu
index f8d58645a1..4f0bd0ceda 100644
--- a/cpp/src/routing/utilities/check_input.cu
+++ b/cpp/src/routing/utilities/check_input.cu
@@ -8,6 +8,7 @@
 #include <cuopt/error.hpp>
 #include <routing/utilities/check_input.hpp>
 
+#include <cuda/stream>
 #include <raft/core/device_mdarray.hpp>
 #include <raft/linalg/transpose.cuh>
 #include <raft/util/cudart_utils.hpp>
@@ -33,7 +34,7 @@ namespace detail {
  * @param stream_view Stream view
  */
 template <typename T>
-void transform_absolute(rmm::device_uvector<T>& v, rmm::cuda_stream_view stream_view)
+void transform_absolute(rmm::device_uvector<T>& v, cuda::stream_ref stream_view)
 {
   thrust::transform(
     rmm::exec_policy(stream_view), v.begin(), v.end(), v.begin(), [] __device__(T x) -> T {
@@ -57,7 +58,7 @@ bool check_pickup_tw(const i_t* pickup_indices,
                      const i_t* earliest_time,
                      const i_t* latest_time,
                      size_t n_requests,
-                     rmm::cuda_stream_view stream_view)
+                     cuda::stream_ref stream_view)
 {
   typedef typename rmm::device_uvector<const i_t>::iterator IterConstInt;
   thrust::permutation_iterator<IterConstInt, IterConstInt> pickup_iter(earliest_time,
@@ -87,7 +88,7 @@ bool check_pickup_demands(const i_t* pickup_indices,
                           const i_t* delivery_indices,
                           const i_t* demands,
                           size_t n_requests,
-                          rmm::cuda_stream_view stream_view)
+                          cuda::stream_ref stream_view)
 {
   typedef typename rmm::device_uvector<const i_t>::iterator IterConstInt;
   thrust::permutation_iterator<IterConstInt, IterConstInt> pickup_iter(demands, pickup_indices);
@@ -107,7 +108,7 @@ bool check_pdp_values(const i_t* pickup_indices,
                       const i_t* delivery_indices,
                       const v_t* values,
                       size_t n_requests,
-                      rmm::cuda_stream_view stream_view)
+                      cuda::stream_ref stream_view)
 {
   auto pickup_iter     = thrust::make_permutation_iterator(values, pickup_indices);
   auto delivery_iter   = thrust::make_permutation_iterator(values, delivery_indices);
@@ -158,7 +159,7 @@ template bool is_symmetric_matrix<int, float>(float const*, int, raft::handle_t
 template <typename i_t>
 bool check_min_latest_with_depot(rmm::device_uvector<i_t>& v_latest_time,
                                  i_t depot_earliest,
-                                 rmm::cuda_stream_view stream_view)
+                                 cuda::stream_ref stream_view)
 {
   i_t min_latest;
   i_t* min_latest_ptr = thrust::min_element(
@@ -178,7 +179,7 @@ bool check_min_latest_with_depot(rmm::device_uvector<i_t>& v_latest_time,
 template <typename i_t>
 bool check_max_earliest_with_depot(rmm::device_uvector<i_t>& v_earliest_time,
                                    i_t depot_latest,
-                                   rmm::cuda_stream_view stream_view)
+                                   cuda::stream_ref stream_view)
 {
   i_t max_earliest;
   i_t* max_earliest_ptr = thrust::max_element(
@@ -198,7 +199,7 @@ bool check_max_earliest_with_depot(rmm::device_uvector<i_t>& v_earliest_time,
 template <typename i_t>
 bool check_earliest_with_latest(rmm::device_uvector<i_t>& v_earliest_time,
                                 rmm::device_uvector<i_t>& v_latest_time,
-                                rmm::cuda_stream_view stream_view)
+                                cuda::stream_ref stream_view)
 {
   return thrust::equal(rmm::exec_policy(stream_view),
                        v_earliest_time.begin(),
@@ -220,7 +221,7 @@ bool check_min_max_values(const T* ptr,
                           size_t size,
                           const RefType min_value,
                           const RefType max_value,
-                          rmm::cuda_stream_view stream_view)
+                          cuda::stream_ref stream_view)
 {
   T min, max;
   thrust::pair<const T*, const T*> pair =
@@ -241,7 +242,7 @@ void check_guess(i_t const* guess_id,
                  i_t fleet_size,
                  bool const* drop_return_trip,
                  bool const* skip_first_trip,
-                 rmm::cuda_stream_view stream_view)
+                 cuda::stream_ref stream_view)
 {
   cuopt_expects(check_min_max_values(truck_id, size, 0, fleet_size - 1, stream_view),
                 error_type_t::ValidationError,
@@ -341,7 +342,7 @@ bool check_no_circular_precedence(i_t node_id,
                                   i_t const* preceding_nodes,
                                   i_t n_preceding_nodes,
                                   std::unordered_map<i_t, std::pair<i_t const*, i_t>> precedence,
-                                  rmm::cuda_stream_view stream_view)
+                                  cuda::stream_ref stream_view)
 {
   for (const auto& pair : precedence) {
     auto other_node   = pair.first;
@@ -372,7 +373,7 @@ bool check_no_circular_precedence(i_t node_id,
  * @return bool Whether the item exists
  */
 template <typename T>
-bool check_exists(T item_id, T const* device_ptr, T n_items, rmm::cuda_stream_view stream_view)
+bool check_exists(T item_id, T const* device_ptr, T n_items, cuda::stream_ref stream_view)
 {
   auto end_ptr  = device_ptr + n_items;
   auto iter_end = thrust::find(rmm::exec_policy(stream_view), device_ptr, end_ptr, item_id);
@@ -383,74 +384,69 @@ template bool check_min_max_values<uint8_t, int>(const uint8_t* ptr,
                                                  size_t size,
                                                  const int min_value,
                                                  const int max_value,
-                                                 rmm::cuda_stream_view stream_view);
+                                                 cuda::stream_ref stream_view);
 
 template bool check_min_max_values<int, int>(const int* ptr,
                                              size_t size,
                                              const int min_value,
                                              const int max_value,
-                                             rmm::cuda_stream_view stream_view);
+                                             cuda::stream_ref stream_view);
 
 template bool check_min_max_values<int, int16_t>(const int* ptr,
                                                  size_t size,
                                                  const int16_t min_value,
                                                  const int16_t max_value,
-                                                 rmm::cuda_stream_view stream_view);
+                                                 cuda::stream_ref stream_view);
 
 template bool check_min_max_values<int, uint16_t>(const int* ptr,
                                                   size_t size,
                                                   const uint16_t min_value,
                                                   const uint16_t max_value,
-                                                  rmm::cuda_stream_view stream_view);
+                                                  cuda::stream_ref stream_view);
 
 template bool check_min_max_values<float, float>(const float* ptr,
                                                  size_t size,
                                                  const float min_value,
                                                  const float max_value,
-                                                 rmm::cuda_stream_view stream_view);
+                                                 cuda::stream_ref stream_view);
 
 template bool check_min_max_values<double, double>(const double* ptr,
                                                    size_t size,
                                                    const double min_value,
                                                    const double max_value,
-                                                   rmm::cuda_stream_view stream_view);
+                                                   cuda::stream_ref stream_view);
 
-template void transform_absolute<int>(rmm::device_uvector<int>& v,
-                                      rmm::cuda_stream_view stream_view);
+template void transform_absolute<int>(rmm::device_uvector<int>& v, cuda::stream_ref stream_view);
 
 template bool check_no_circular_precedence<int>(
   int node_id,
   int const* preceding_nodes,
   int n_preceding_nodes,
   std::unordered_map<int, std::pair<int const*, int>> precedence,
-  rmm::cuda_stream_view stream_view);
+  cuda::stream_ref stream_view);
 
 template bool check_exists<int>(int item_id,
                                 int const* device_ptr,
                                 int n_items,
-                                rmm::cuda_stream_view stream_view);
+                                cuda::stream_ref stream_view);
 
 template bool check_earliest_with_latest<int>(rmm::device_uvector<int>&,
                                               rmm::device_uvector<int>&,
-                                              rmm::cuda_stream_view);
-template bool check_max_earliest_with_depot<int>(rmm::device_uvector<int>&,
-                                                 int,
-                                                 rmm::cuda_stream_view);
+                                              cuda::stream_ref);
+template bool check_max_earliest_with_depot<int>(rmm::device_uvector<int>&, int, cuda::stream_ref);
 template bool check_pickup_tw<int>(
-  int const*, int const*, int const*, int const*, unsigned long, rmm::cuda_stream_view);
+  int const*, int const*, int const*, int const*, unsigned long, cuda::stream_ref);
 template bool check_pickup_demands<int>(
-  int const*, int const*, int const*, unsigned long, rmm::cuda_stream_view);
+  int const*, int const*, int const*, unsigned long, cuda::stream_ref);
 
 template bool check_pdp_values<int, uint8_t>(
-  int const*, int const*, uint8_t const*, unsigned long, rmm::cuda_stream_view);
+  int const*, int const*, uint8_t const*, unsigned long, cuda::stream_ref);
 template bool check_pdp_values<int, int>(
-  int const*, int const*, int const*, unsigned long, rmm::cuda_stream_view);
+  int const*, int const*, int const*, unsigned long, cuda::stream_ref);
 template bool check_pdp_values<int, float>(
-  int const*, int const*, float const*, unsigned long, rmm::cuda_stream_view);
+  int const*, int const*, float const*, unsigned long, cuda::stream_ref);
 
-template bool check_min_latest_with_depot<int>(rmm::device_uvector<int>&,
-                                               int,
-                                               rmm::cuda_stream_view);
+template bool check_min_latest_with_depot<int>(rmm::device_uvector<int>&, int, cuda::stream_ref);
 template void check_guess<int>(int const*,
                                int const*,
                                int const*,
@@ -460,7 +456,7 @@ template void check_guess<int>(int const*,
                                int,
                                bool const*,
                                bool const*,
-                               rmm::cuda_stream_view);
+                               cuda::stream_ref);
 
 }  // namespace detail
 }  // namespace routing
diff --git a/cpp/src/routing/utilities/check_input.hpp b/cpp/src/routing/utilities/check_input.hpp
index 07ce697c1a..449df036f1 100644
--- a/cpp/src/routing/utilities/check_input.hpp
+++ b/cpp/src/routing/utilities/check_input.hpp
@@ -1,12 +1,13 @@
 /* clang-format off */
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 /* clang-format on */
 
 #pragma once
 
+#include <cuda/stream>
 #include <raft/core/handle.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
@@ -23,21 +24,21 @@ bool is_symmetric_matrix(f_t const* matrix, i_t width, raft::handle_t const* han
 template <typename i_t>
 bool check_min_latest_with_depot(rmm::device_uvector<i_t>& v_latest_time,
                                  i_t depot_earliest,
-                                 rmm::cuda_stream_view stream_view);
+                                 cuda::stream_ref stream_view);
 template <typename i_t>
 bool check_max_earliest_with_depot(rmm::device_uvector<i_t>& v_earliest_time,
                                    i_t depot_latest,
-                                   rmm::cuda_stream_view stream_view);
+                                   cuda::stream_ref stream_view);
 template <typename i_t>
 bool check_earliest_with_latest(rmm::device_uvector<i_t>& v_earliest_time,
                                 rmm::device_uvector<i_t>& v_latest_time,
-                                rmm::cuda_stream_view stream_view);
+                                cuda::stream_ref stream_view);
 template <typename T, typename RefType>
 bool check_min_max_values(const T* ptr,
                           size_t size,
                           const RefType min_value,
                           const RefType max_value,
-                          rmm::cuda_stream_view stream_view);
+                          cuda::stream_ref stream_view);
 
 template <typename i_t>
 void check_guess(i_t const* guess_id,
@@ -49,7 +50,7 @@ void check_guess(i_t const* guess_id,
                  i_t fleet_size,
                  bool const* drop_return_trip,
                  bool const* skip_first_trip,
-                 rmm::cuda_stream_view stream_view);
+                 cuda::stream_ref stream_view);
 
 template <typename i_t>
 bool check_pickup_tw(const i_t* pickup_indices,
@@ -57,34 +58,34 @@ bool check_pickup_tw(const i_t* pickup_indices,
                      const i_t* earliest_time,
                      const i_t* latest_time,
                      size_t n_requests,
-                     rmm::cuda_stream_view stream_view);
+                     cuda::stream_ref stream_view);
 
 template <typename i_t>
 bool check_pickup_demands(const i_t* pickup_indices,
                           const i_t* delivery_indices,
                           const i_t* demands,
                           size_t n_requests,
-                          rmm::cuda_stream_view stream_view);
+                          cuda::stream_ref stream_view);
 
 template <typename i_t, typename v_t>
 bool check_pdp_values(const i_t* pickup_indices,
                       const i_t* delivery_indices,
                       const v_t* values,
                       size_t n_requests,
-                      rmm::cuda_stream_view stream_view);
+                      cuda::stream_ref stream_view);
 
 template <typename i_t>
 bool check_no_circular_precedence(i_t node_id,
                                   i_t const* preceding_nodes,
                                   i_t n_preceding_nodes,
                                   std::unordered_map<i_t, std::pair<i_t const*, i_t>> precedence,
-                                  rmm::cuda_stream_view stream_view);
+                                  cuda::stream_ref stream_view);
 
 template <typename T>
-bool check_exists(T item_id, T const* device_ptr, T n_items, rmm::cuda_stream_view stream_view);
+bool check_exists(T item_id, T const* device_ptr, T n_items, cuda::stream_ref stream_view);
 
 template <typename T>
-void transform_absolute(rmm::device_uvector<T>& v, rmm::cuda_stream_view stream_view);
+void transform_absolute(rmm::device_uvector<T>& v, cuda::stream_ref stream_view);
 
 }  // namespace detail
 }  // namespace routing
diff --git a/cpp/src/routing/utilities/md_utils.hpp b/cpp/src/routing/utilities/md_utils.hpp
index 7de1be3f0a..4edeb9bcb3 100644
--- a/cpp/src/routing/utilities/md_utils.hpp
+++ b/cpp/src/routing/utilities/md_utils.hpp
@@ -1,6 +1,6 @@
 /* clang-format off */
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 /* clang-format on */
@@ -12,6 +12,7 @@
 
 #include <thrust/functional.h>
 #include <thrust/logical.h>
+#include <cuda/stream>
 #include <rmm/device_uvector.hpp>
 #include <utilities/copy_helpers.hpp>
 #include <utilities/macros.cuh>
@@ -95,8 +96,8 @@ struct h_mdarray_t {
 
 template <typename f_t, size_t NCON_DIMS = 4>
 struct d_mdarray_t {
-  d_mdarray_t(rmm::cuda_stream_view stream_) : buffer(0, stream_), stream(stream_) {}
-  d_mdarray_t(std::vector<size_t> const& extent_, rmm::cuda_stream_view stream_)
+  d_mdarray_t(cuda::stream_ref stream_) : buffer(0, stream_), stream(stream_) {}
+  d_mdarray_t(std::vector<size_t> const& extent_, cuda::stream_ref stream_)
     : buffer(0, stream_), stream(stream_)
   {
     cuopt_assert(extent_.size() == NCON_DIMS, "Wrong dimensions");
@@ -138,7 +139,7 @@ struct d_mdarray_t {
 
   size_t extent[NCON_DIMS];
   rmm::device_uvector<f_t> buffer;
-  rmm::cuda_stream_view stream;
+  cuda::stream_ref stream;
 };
 
 namespace detail {
@@ -196,7 +197,7 @@ template <typename f_t>
 auto create_device_mdarray(size_t nlocations,
                            uint8_t n_vehicle_types,
                            uint8_t n_matrix_types,
-                           rmm::cuda_stream_view stream)
+                           cuda::stream_ref stream)
 {
   std::vector<size_t> full_matrix_extent{n_vehicle_types, n_matrix_types, nlocations, nlocations};
   d_mdarray_t<f_t> matrices{full_matrix_extent, stream};
@@ -204,7 +205,7 @@ auto create_device_mdarray(size_t nlocations,
 }
 
 inline auto get_unique_vehicle_types(const raft::device_span<uint8_t const>& vehicle_types,
-                                     rmm::cuda_stream_view stream)
+                                     cuda::stream_ref stream)
 {
   auto h_vehicle_types = cuopt::host_copy(vehicle_types, stream);
 
diff --git a/cpp/src/utilities/copy_helpers.hpp b/cpp/src/utilities/copy_helpers.hpp
index 211fd4552a..4c228f0393 100644
--- a/cpp/src/utilities/copy_helpers.hpp
+++ b/cpp/src/utilities/copy_helpers.hpp
@@ -7,6 +7,7 @@
 
 #pragma once
 
+#include <cuda/stream>
 #include <raft/core/device_span.hpp>
 #include <raft/core/handle.hpp>
 #include <raft/util/cudart_utils.hpp>
@@ -119,7 +120,7 @@ __host__ __device__ inline typename scalar_type<f_t2>::type& get_upper(f_t2& val
  * @return auto
  */
 template <typename T>
-auto host_copy(T const* device_ptr, size_t size, rmm::cuda_stream_view stream_view)
+auto host_copy(T const* device_ptr, size_t size, cuda::stream_ref stream_view)
 {
   if (!device_ptr) return std::vector<T>{};
   std::vector<T> host_vec(size);
@@ -137,7 +138,7 @@ auto host_copy(T const* device_ptr, size_t size, rmm::cuda_stream_view stream_vi
  * @param[in] stream_view
  * @return auto
  */
-inline auto host_copy(bool const* device_ptr, size_t size, rmm::cuda_stream_view stream_view)
+inline auto host_copy(bool const* device_ptr, size_t size, cuda::stream_ref stream_view)
 {
   if (!device_ptr) { return std::vector<bool>(0); }
   rmm::device_uvector<int> d_int_vec(size, stream_view);
@@ -163,7 +164,7 @@ inline auto host_copy(bool const* device_ptr, size_t size, rmm::cuda_stream_view
  * @return auto
  */
 template <typename T, typename Allocator>
-auto host_copy(rmm::device_uvector<T> const& device_vec, rmm::cuda_stream_view stream_view)
+auto host_copy(rmm::device_uvector<T> const& device_vec, cuda::stream_ref stream_view)
 {
   std::vector<T, Allocator> host_vec(device_vec.size());
   raft::copy(host_vec.data(), device_vec.data(), device_vec.size(), stream_view);
@@ -180,7 +181,7 @@ auto host_copy(rmm::device_uvector<T> const& device_vec, rmm::cuda_stream_view s
  * @return auto
  */
 template <typename T>
-auto host_copy(raft::device_span<T> const& device_vec, rmm::cuda_stream_view stream_view)
+auto host_copy(raft::device_span<T> const& device_vec, cuda::stream_ref stream_view)
 {
   return host_copy(device_vec.data(), device_vec.size(), stream_view);
 }
@@ -194,7 +195,7 @@ auto host_copy(raft::device_span<T> const& device_vec, rmm::cuda_stream_view str
  * @return auto
  */
 template <typename T>
-auto host_copy(rmm::device_uvector<T> const& device_vec, rmm::cuda_stream_view stream_view)
+auto host_copy(rmm::device_uvector<T> const& device_vec, cuda::stream_ref stream_view)
 {
   return host_copy(device_vec.data(), device_vec.size(), stream_view);
 }
@@ -209,7 +210,7 @@ auto host_copy(rmm::device_uvector<T> const& device_vec, rmm::cuda_stream_view s
  */
 template <typename T>
 inline rmm::device_uvector<T> device_copy(rmm::device_uvector<T> const& device_vec,
-                                          rmm::cuda_stream_view stream_view)
+                                          cuda::stream_ref stream_view)
 {
   rmm::device_uvector<T> device_vec_copy(device_vec.size(), stream_view);
   raft::copy(device_vec_copy.data(), device_vec.data(), device_vec.size(), stream_view);
@@ -227,7 +228,7 @@ inline rmm::device_uvector<T> device_copy(rmm::device_uvector<T> const& device_v
 template <typename T>
 inline void device_copy(rmm::device_uvector<T>& device_vec,
                         std::vector<T> const& host_vec,
-                        rmm::cuda_stream_view stream_view)
+                        cuda::stream_ref stream_view)
 {
   device_vec.resize(host_vec.size(), stream_view);
   raft::copy(device_vec.data(), host_vec.data(), host_vec.size(), stream_view);
@@ -242,8 +243,7 @@ inline void device_copy(rmm::device_uvector<T>& device_vec,
  * @return device_vec
  */
 template <typename T, typename Allocator>
-inline auto device_copy(std::vector<T, Allocator> const& host_vec,
-                        rmm::cuda_stream_view stream_view)
+inline auto device_copy(std::vector<T, Allocator> const& host_vec, cuda::stream_ref stream_view)
 {
   rmm::device_uvector<T> device_vec(host_vec.size(), stream_view);
   raft::copy(device_vec.data(), host_vec.data(), host_vec.size(), stream_view);
@@ -257,7 +257,7 @@ inline auto device_copy(std::vector<T, Allocator> const& host_vec,
  * @param[in] stream_view
  * @return device_vec
  */
-inline auto device_copy(std::vector<bool> const& host_vec, rmm::cuda_stream_view stream_view)
+inline auto device_copy(std::vector<bool> const& host_vec, cuda::stream_ref stream_view)
 {
   std::vector<uint8_t> host_vec_int(host_vec.size());
   for (size_t i = 0; i < host_vec.size(); ++i) {
@@ -340,7 +340,7 @@ raft::device_span<const T> make_span(rmm::device_uvector<T> const& container)
 template <typename T>
 inline void expand_device_copy(rmm::device_uvector<T>& device_vec,
                                std::vector<T> const& host_vec,
-                               rmm::cuda_stream_view stream_view)
+                               cuda::stream_ref stream_view)
 {
   if (host_vec.size() > device_vec.size()) { device_vec.resize(host_vec.size(), stream_view); }
   raft::copy(device_vec.data(), host_vec.data(), host_vec.size(), stream_view);
@@ -349,7 +349,7 @@ inline void expand_device_copy(rmm::device_uvector<T>& device_vec,
 template <typename T>
 inline void expand_device_copy(rmm::device_uvector<T>& dst_vec,
                                rmm::device_uvector<T> const& src_vec,
-                               rmm::cuda_stream_view stream_view)
+                               cuda::stream_ref stream_view)
 {
   if (src_vec.size() > dst_vec.size()) { dst_vec.resize(src_vec.size(), stream_view); }
   raft::copy(dst_vec.data(), src_vec.data(), src_vec.size(), stream_view);
diff --git a/cpp/src/utilities/event_handler.cuh b/cpp/src/utilities/event_handler.cuh
index 17f6bd787c..590e4687cc 100644
--- a/cpp/src/utilities/event_handler.cuh
+++ b/cpp/src/utilities/event_handler.cuh
@@ -7,8 +7,8 @@
 #pragma once
 
 #include <cuda_runtime.h>
+#include <cuda/stream>
 #include <raft/util/cuda_utils.cuh>
-#include <rmm/cuda_stream_view.hpp>
 
 namespace cuopt {
 
@@ -21,17 +21,17 @@ class event_handler_t {
   event_handler_t(const event_handler_t&)            = delete;
   event_handler_t& operator=(const event_handler_t&) = delete;
 
-  void record(rmm::cuda_stream_view stream_view)
+  void record(cuda::stream_ref stream_view)
   {
     RAFT_CUDA_TRY(cudaEventRecord(event_, stream_view.get()));
   }
 
-  void record_with_flags(rmm::cuda_stream_view stream_view, int flags)
+  void record_with_flags(cuda::stream_ref stream_view, int flags)
   {
     RAFT_CUDA_TRY(cudaEventRecordWithFlags(event_, stream_view.get(), flags));
   }
 
-  void stream_wait(rmm::cuda_stream_view stream_view)
+  void stream_wait(cuda::stream_ref stream_view)
   {
     RAFT_CUDA_TRY(cudaStreamWaitEvent(stream_view.get(), event_));
   }
diff --git a/cpp/src/utilities/manual_cuda_graph.cuh b/cpp/src/utilities/manual_cuda_graph.cuh
index bdc5ba9fd4..6c892b915a 100644
--- a/cpp/src/utilities/manual_cuda_graph.cuh
+++ b/cpp/src/utilities/manual_cuda_graph.cuh
@@ -10,9 +10,9 @@
 #include <utilities/logger.hpp>
 #include <utilities/macros.cuh>
 
+#include <cuda/stream>
 #include <raft/core/error.hpp>
 #include <raft/util/cudart_utils.hpp>
-#include <rmm/cuda_stream_view.hpp>
 
 #include <cuda_runtime.h>
 
@@ -68,7 +68,7 @@ class manual_cuda_graph_t {
   ~manual_cuda_graph_t() { destroy(); }
 
   template <typename F>
-  void run(rmm::cuda_stream_view stream, F&& work)
+  void run(cuda::stream_ref stream, F&& work)
   {
     if (instance_ != nullptr) {
       RAFT_CUDA_TRY(cudaGraphLaunch(instance_, stream.get()));
diff --git a/cpp/src/utilities/vector_helpers.cuh b/cpp/src/utilities/vector_helpers.cuh
index 91f35c34fc..e03460529d 100644
--- a/cpp/src/utilities/vector_helpers.cuh
+++ b/cpp/src/utilities/vector_helpers.cuh
@@ -8,6 +8,7 @@
 #pragma once
 
 #include <thrust/extrema.h>
+#include <cuda/stream>
 #include <raft/core/device_span.hpp>
 #include <raft/core/handle.hpp>
 #include <rmm/device_uvector.hpp>
@@ -39,7 +40,7 @@ __global__ void sequence_with_multiplier_kernel(T* data_ptr, int mult, size_t si
 }
 
 template <typename T>
-void async_fill(rmm::device_uvector<T>& vec, T item, rmm::cuda_stream_view stream)
+void async_fill(rmm::device_uvector<T>& vec, T item, cuda::stream_ref stream)
 {
   constexpr size_t TPB = 256;
   size_t n_blocks      = (vec.size() + TPB - 1) / TPB;
@@ -47,7 +48,7 @@ void async_fill(rmm::device_uvector<T>& vec, T item, rmm::cuda_stream_view strea
 }
 
 template <typename T>
-void async_fill(T* vec, T item, size_t size, rmm::cuda_stream_view stream)
+void async_fill(T* vec, T item, size_t size, cuda::stream_ref stream)
 {
   constexpr size_t TPB = 256;
   size_t n_blocks      = (size + TPB - 1) / TPB;
@@ -55,7 +56,7 @@ void async_fill(T* vec, T item, size_t size, rmm::cuda_stream_view stream)
 }
 
 template <typename T>
-void async_sequence(rmm::device_uvector<T>& vec, rmm::cuda_stream_view stream)
+void async_sequence(rmm::device_uvector<T>& vec, cuda::stream_ref stream)
 {
   constexpr size_t TPB = 256;
   size_t n_blocks      = (vec.size() + TPB - 1) / TPB;
@@ -63,9 +64,7 @@ void async_sequence(rmm::device_uvector<T>& vec, rmm::cuda_stream_view stream)
 }
 
 template <typename T>
-void async_sequence_with_multiplier(rmm::device_uvector<T>& vec,
-                                    int mult,
-                                    rmm::cuda_stream_view stream)
+void async_sequence_with_multiplier(rmm::device_uvector<T>& vec, int mult, cuda::stream_ref stream)
 {
   constexpr size_t TPB = 256;
   size_t n_blocks      = (vec.size() + TPB - 1) / TPB;
diff --git a/cpp/tests/linear_programming/utilities/pdlp_test_utilities.cuh b/cpp/tests/linear_programming/utilities/pdlp_test_utilities.cuh
index 7e3f83dae9..cff05578fe 100644
--- a/cpp/tests/linear_programming/utilities/pdlp_test_utilities.cuh
+++ b/cpp/tests/linear_programming/utilities/pdlp_test_utilities.cuh
@@ -6,6 +6,7 @@
 /* clang-format on */
 #pragma once
 
+#include <cuda/stream>
 #include <cuopt/mathematical_optimization/optimization_problem.hpp>
 #include <cuopt/mathematical_optimization/pdlp/solver_settings.hpp>
 #include <cuopt/mathematical_optimization/pdlp/solver_solution.hpp>
@@ -53,7 +54,7 @@ static cuopt::mathematical_optimization::optimization_problem_solution_t<i_t, f_
 template <typename f_t>
 static void assign_device_uvector_from_host(rmm::device_uvector<f_t>& target,
                                             const std::vector<f_t>& src,
-                                            rmm::cuda_stream_view stream)
+                                            cuda::stream_ref stream)
 {
   target.resize(src.size(), stream);
   raft::copy(target.data(), src.data(), src.size(), stream);
diff --git a/cpp/tests/routing/routing_test.cuh b/cpp/tests/routing/routing_test.cuh
index cdafbbf1f7..839372db21 100644
--- a/cpp/tests/routing/routing_test.cuh
+++ b/cpp/tests/routing/routing_test.cuh
@@ -21,6 +21,7 @@
 #include <utilities/copy_helpers.hpp>
 #include <utilities/high_res_timer.hpp>
 
+#include <cuda/stream>
 #include <raft/core/error.hpp>
 #include <raft/core/handle.hpp>
 
@@ -689,7 +690,7 @@ class base_test_t {
   }
 
   raft::handle_t handle_;
-  rmm::cuda_stream_view stream_view_;
+  cuda::stream_ref stream_view_;
   HighResTimer hr_timer_;
 
   bool multi_capacity_{false};
diff --git a/cpp/tests/socp/sparse_augmented_kkt_test.cu b/cpp/tests/socp/sparse_augmented_kkt_test.cu
index 533b29cceb..1ddc3c8f7e 100644
--- a/cpp/tests/socp/sparse_augmented_kkt_test.cu
+++ b/cpp/tests/socp/sparse_augmented_kkt_test.cu
@@ -25,8 +25,7 @@ namespace cuopt::mathematical_optimization::barrier::test {
 namespace {
 
 // Packed Hs_diag reference: eta^2 on every entry, head scaled by rank-2 corner d.
-std::vector<double> expected_Hs_diag(const cone_data_t<int, double>& cones,
-                                     rmm::cuda_stream_view stream)
+std::vector<double> expected_Hs_diag(const cone_data_t<int, double>& cones, cuda::stream_ref stream)
 {
   const int E                    = static_cast<int>(cones.n_sparse_cone_entries);
   auto d_host                    = cuopt::host_copy(cones.d, stream);
diff --git a/skills/cuopt-developer/benchmark/evals.json b/skills/cuopt-developer/benchmark/evals.json
index d1a0b4f1d0..e845fe6c98 100644
--- a/skills/cuopt-developer/benchmark/evals.json
+++ b/skills/cuopt-developer/benchmark/evals.json
@@ -50,11 +50,11 @@
     "question": "I'm writing a new CUDA kernel in cpp/src for cuOpt. How should I allocate device buffers for it?",
     "expected_skill": "cuopt-developer",
     "expected_script": null,
-    "ground_truth": "The agent prescribes RMM (rmm::device_uvector or another RMM allocator) for all device-side allocations and explicitly forbids raw new/delete or cudaMalloc. It explains that operations should accept and use a cuda_stream_view so allocations and kernel launches are stream-ordered, points to existing kernels in cpp/src as reference for RMM allocation, RAFT utilities, and kernel launch patterns, and mentions RAFT_CUDA_TRY for CUDA error checking.",
+    "ground_truth": "The agent prescribes RMM (rmm::device_uvector or another RMM allocator) for all device-side allocations and explicitly forbids raw new/delete or cudaMalloc. It explains that operations should accept and use a cuda::stream_ref so allocations and kernel launches are stream-ordered, points to existing kernels in cpp/src as reference for RMM allocation, RAFT utilities, and kernel launch patterns, and mentions RAFT_CUDA_TRY for CUDA error checking.",
     "expected_behavior": [
       "Recommends rmm::device_uvector or another RMM allocator",
       "Explicitly says raw new/delete and cudaMalloc are not allowed",
-      "Mentions cuda_stream_view and stream-ordered operations",
+      "Mentions cuda::stream_ref and stream-ordered operations",
       "Points the user to existing code in cpp/src as reference for patterns",
       "Mentions RAFT utilities or RAFT_CUDA_TRY for CUDA error checking"
     ]
@@ -461,13 +461,13 @@
     "question": "Does cuOpt use RAFT or RMM? What conventions should I follow when writing GPU code in the codebase?",
     "expected_skill": "cuopt-developer",
     "expected_script": null,
-    "ground_truth": "cuOpt uses both. RMM provides device-memory allocators (rmm::device_uvector and similar); raw new/delete or cudaMalloc are not allowed. RAFT provides utilities including RAFT_CUDA_TRY for wrapping CUDA API calls so failures throw with context. Operations are stream-ordered via cuda_stream_view; views (the _view suffix) are non-owning. The agent points to existing code in cpp/src/ as reference for these patterns.",
+    "ground_truth": "cuOpt uses both. RMM provides device-memory allocators (rmm::device_uvector and similar); raw new/delete or cudaMalloc are not allowed. RAFT provides utilities including RAFT_CUDA_TRY for wrapping CUDA API calls so failures throw with context. Operations are stream-ordered via cuda::stream_ref, which is non-owning. The agent points to existing code in cpp/src/ as reference for these patterns.",
     "expected_behavior": [
       "States cuOpt uses both RAFT and RMM",
       "Mentions rmm::device_uvector (or RMM allocators) for device memory",
       "Mentions RAFT_CUDA_TRY for CUDA error wrapping",
-      "Mentions cuda_stream_view and stream-ordered operations",
-      "Mentions _view suffix means non-owning",
+      "Mentions cuda::stream_ref and stream-ordered operations",
+      "Mentions cuda::stream_ref is non-owning",
       "Points to existing cpp/src/ code as the reference for patterns"
     ]
   },
diff --git a/skills/cuopt-developer/references/conventions.md b/skills/cuopt-developer/references/conventions.md
index 1bef2bbe3a..5c2ea9da80 100644
--- a/skills/cuopt-developer/references/conventions.md
+++ b/skills/cuopt-developer/references/conventions.md
@@ -192,8 +192,8 @@ int* data = new int[100];
 rmm::device_uvector<int> data(100, stream);
 ```
 
-- All operations should accept `cuda_stream_view`
-- Views (`*_view` suffix) are non-owning
+- All operations should accept `cuda::stream_ref`
+- Stream references are non-owning
 
 Read existing code in `cpp/src/` for real examples of RMM allocation, stream-ordering, RAFT utilities, and kernel launch patterns.