diff --git a/cpp/benchmarks/cuda_stream_pool/cuda_stream_pool_bench.cpp b/cpp/benchmarks/cuda_stream_pool/cuda_stream_pool_bench.cpp
index fba7f2664..096511964 100644
--- a/cpp/benchmarks/cuda_stream_pool/cuda_stream_pool_bench.cpp
+++ b/cpp/benchmarks/cuda_stream_pool/cuda_stream_pool_bench.cpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -18,7 +18,7 @@ static void BM_StreamPoolGetStream(benchmark::State& state)
 
   for (auto _ : state) {  // NOLINT(clang-analyzer-deadcode.DeadStores)
     auto stream = stream_pool.get_stream();
-    cudaStreamQuery(stream.value());
+    cudaStreamQuery(cuda::stream_ref{stream}.get());
   }
 
   state.SetItemsProcessed(static_cast<int64_t>(state.iterations()));
@@ -29,7 +29,7 @@ static void BM_CudaStreamClass(benchmark::State& state)
 {
   for (auto _ : state) {  // NOLINT(clang-analyzer-deadcode.DeadStores)
     auto stream = rmm::cuda_stream{};
-    cudaStreamQuery(stream.view().value());
+    cudaStreamQuery(cuda::stream_ref{stream}.get());
   }
 
   state.SetItemsProcessed(static_cast<int64_t>(state.iterations()));
diff --git a/cpp/benchmarks/device_uvector/device_uvector_bench.cu b/cpp/benchmarks/device_uvector/device_uvector_bench.cu
index b28b2ff7a..33fe0c80b 100644
--- a/cpp/benchmarks/device_uvector/device_uvector_bench.cu
+++ b/cpp/benchmarks/device_uvector/device_uvector_bench.cu
@@ -15,6 +15,7 @@
 #include <rmm/mr/per_device_resource.hpp>
 #include <rmm/mr/pool_memory_resource.hpp>
 
+#include <cuda/stream_ref>
 #include <cuda_runtime_api.h>
 #include <thrust/device_vector.h>
 #include <thrust/memory.h>
@@ -31,7 +32,7 @@ void BM_UvectorSizeConstruction(benchmark::State& state)
 
   for (auto _ : state) {  // NOLINT(clang-analyzer-deadcode.DeadStores)
     rmm::device_uvector<std::int32_t> vec(static_cast<std::size_t>(state.range(0)),
-                                          rmm::cuda_stream_view{});
+                                          cuda::stream_ref{cudaStream_t{nullptr}});
     cudaDeviceSynchronize();
   }
 
@@ -78,7 +79,7 @@ using rmm_vector    = rmm::device_vector<int32_t>;
 using rmm_uvector   = rmm::device_uvector<int32_t>;
 
 template <typename Vector>
-Vector make_vector(std::size_t num_elements, rmm::cuda_stream_view stream, bool zero_init = false)
+Vector make_vector(std::size_t num_elements, cuda::stream_ref stream, bool zero_init = false)
 {
   static_assert(std::is_same_v<Vector, thrust_vector> or std::is_same_v<Vector, rmm_vector> or
                   std::is_same_v<Vector, rmm_uvector>,
@@ -90,7 +91,7 @@ Vector make_vector(std::size_t num_elements, rmm::cuda_stream_view stream, bool
   } else if constexpr (std::is_same_v<Vector, rmm_uvector>) {
     auto vec = Vector(num_elements, stream);
     if (zero_init) {
-      cudaMemsetAsync(vec.data(), 0, num_elements * sizeof(std::int32_t), stream.value());
+      cudaMemsetAsync(vec.data(), 0, num_elements * sizeof(std::int32_t), stream.get());
     }
     return vec;
   }
@@ -111,14 +112,14 @@ void vector_workflow(std::size_t num_elements,
 {
   auto input = make_vector<Vector>(num_elements, input_stream, true);
   input_stream.synchronize();
-  for (rmm::cuda_stream_view stream : streams) {
+  for (cuda::stream_ref stream : streams) {
     auto output = make_vector<Vector>(num_elements, stream);
-    kernel<<<num_blocks, block_size, 0, stream.value()>>>(
+    kernel<<<num_blocks, block_size, 0, stream.get()>>>(
       vector_data(input), vector_data(output), num_elements);
   }
 
-  for (rmm::cuda_stream_view stream : streams) {
-    stream.synchronize();
+  for (cuda::stream_ref stream : streams) {
+    RMM_CUDA_TRY(cudaStreamSynchronize(stream.get()));
   }
 }
 
diff --git a/cpp/benchmarks/multi_stream_allocations/multi_stream_allocations_bench.cu b/cpp/benchmarks/multi_stream_allocations/multi_stream_allocations_bench.cu
index 093dbf8e9..b578f7e2a 100644
--- a/cpp/benchmarks/multi_stream_allocations/multi_stream_allocations_bench.cu
+++ b/cpp/benchmarks/multi_stream_allocations/multi_stream_allocations_bench.cu
@@ -15,6 +15,7 @@
 #include <rmm/mr/pool_memory_resource.hpp>
 #include <rmm/resource_ref.hpp>
 
+#include <cuda/stream_ref>
 #include <cuda_runtime_api.h>
 
 #include <benchmark/benchmark.h>
@@ -54,9 +55,9 @@ static void run_test(std::size_t num_kernels,
                      rmm::device_async_resource_ref mr)
 {
   for (std::size_t i = 0; i < num_kernels; i++) {
-    auto stream = stream_pool.get_stream(i);
+    auto stream = cuda::stream_ref{stream_pool.get_stream(i)};
     auto buffer = rmm::device_uvector<int64_t>(1, stream, mr);
-    compute_bound_kernel<<<1, 1, 0, stream.value()>>>(buffer.data());
+    compute_bound_kernel<<<1, 1, 0, stream.get()>>>(buffer.data());
   }
 }
 
diff --git a/cpp/benchmarks/random_allocations/random_allocations.cpp b/cpp/benchmarks/random_allocations/random_allocations.cpp
index f342ee07b..16b1f9f6c 100644
--- a/cpp/benchmarks/random_allocations/random_allocations.cpp
+++ b/cpp/benchmarks/random_allocations/random_allocations.cpp
@@ -13,6 +13,9 @@
 #include <rmm/mr/pool_memory_resource.hpp>
 #include <rmm/resource_ref.hpp>
 
+#include <cuda/stream_ref>
+#include <cuda_runtime_api.h>
+
 #include <benchmark/benchmark.h>
 #include <benchmarks/utilities/cxxopts.hpp>
 
@@ -54,7 +57,7 @@ void random_allocation_free(rmm::device_async_resource_ref mr,
                             SizeDistribution size_distribution,
                             std::size_t num_allocations,
                             std::size_t max_usage,  // in MiB
-                            rmm::cuda_stream_view stream = {})
+                            cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}})
 {
   std::default_random_engine generator;
 
@@ -132,7 +135,7 @@ void uniform_random_allocations(
   std::size_t num_allocations,      // NOLINT(bugprone-easily-swappable-parameters)
   std::size_t max_allocation_size,  // size in MiB
   std::size_t max_usage,
-  rmm::cuda_stream_view stream = {})
+  cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}})
 {
   std::uniform_int_distribution<std::size_t> size_distribution(1, max_allocation_size * size_mb);
   random_allocation_free(mr, size_distribution, num_allocations, max_usage, stream);
@@ -144,7 +147,7 @@ void uniform_random_allocations(
                                 std::size_t mean_allocation_size = 500, // in MiB
                                 std::size_t stddev_allocation_size = 500, // in MiB
                                 std::size_t max_usage = 8 << 20,
-                                cuda_stream_view stream) {
+                                cuda::stream_ref stream) {
   std::normal_distribution<std::size_t> size_distribution(, max_allocation_size * size_mb);
 }*/
 
diff --git a/cpp/benchmarks/replay/replay.cpp b/cpp/benchmarks/replay/replay.cpp
index ef6acaaa3..4bdf002c3 100644
--- a/cpp/benchmarks/replay/replay.cpp
+++ b/cpp/benchmarks/replay/replay.cpp
@@ -4,7 +4,6 @@
  */
 
 #include <rmm/aligned.hpp>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/detail/error.hpp>
 #include <rmm/logger.hpp>
 #include <rmm/mr/arena_memory_resource.hpp>
@@ -16,6 +15,7 @@
 #include <rmm/resource_ref.hpp>
 
 #include <cuda/iterator>
+#include <cuda/stream_ref>
 #include <thrust/execution_policy.h>
 #include <thrust/reduce.h>
 
@@ -253,8 +253,14 @@ std::vector<std::vector<rmm::detail::event>> parse_per_thread_events(std::string
                           [](auto const& event) {
                             cudaStream_t custream;
                             memcpy(&custream, &event.stream, sizeof(cudaStream_t));
-                            auto stream = rmm::cuda_stream_view{custream};
-                            return stream.is_default() or stream.is_per_thread_default();
+                            auto stream = cuda::stream_ref{custream};
+#ifdef CUDA_API_PER_THREAD_DEFAULT_STREAM
+                            return stream.get() == cudaStreamLegacy or
+                                   stream.get() == cudaStreamPerThread or stream.get() == nullptr;
+#else
+                            return stream.get() == cudaStreamLegacy or stream.get() == nullptr or
+                                   stream.get() == cudaStreamPerThread;
+#endif
                           }),
               "Non-default streams not currently supported.");
 
diff --git a/cpp/benchmarks/synchronization/synchronization.cpp b/cpp/benchmarks/synchronization/synchronization.cpp
index b33f514f1..2699d085b 100644
--- a/cpp/benchmarks/synchronization/synchronization.cpp
+++ b/cpp/benchmarks/synchronization/synchronization.cpp
@@ -19,7 +19,7 @@
 
 cuda_event_timer::cuda_event_timer(benchmark::State& state,
                                    bool flush_l2_cache,
-                                   rmm::cuda_stream_view stream)
+                                   cuda::stream_ref stream)
   : stream(stream), p_state(&state)
 {
   // flush all of L2$
@@ -36,18 +36,18 @@ cuda_event_timer::cuda_event_timer(benchmark::State& state,
       RMM_CUDA_TRY(cudaMemsetAsync(l2_cache_buffer.data(),
                                    memset_value,
                                    static_cast<std::size_t>(l2_cache_bytes),
-                                   stream.value()));
+                                   stream.get()));
     }
   }
 
   RMM_CUDA_TRY(cudaEventCreate(&start));
   RMM_CUDA_TRY(cudaEventCreate(&stop));
-  RMM_CUDA_TRY(cudaEventRecord(start, stream.value()));
+  RMM_CUDA_TRY(cudaEventRecord(start, stream.get()));
 }
 
 cuda_event_timer::~cuda_event_timer()
 {
-  RMM_CUDA_ASSERT_OK(cudaEventRecord(stop, stream.value()));
+  RMM_CUDA_ASSERT_OK(cudaEventRecord(stop, stream.get()));
   RMM_CUDA_ASSERT_OK(cudaEventSynchronize(stop));
 
   float milliseconds = 0.0F;
diff --git a/cpp/benchmarks/synchronization/synchronization.hpp b/cpp/benchmarks/synchronization/synchronization.hpp
index 556811977..343e3e84e 100644
--- a/cpp/benchmarks/synchronization/synchronization.hpp
+++ b/cpp/benchmarks/synchronization/synchronization.hpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2021, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -47,9 +47,7 @@
 
 #pragma once
 
-#include <rmm/cuda_stream_view.hpp>
-
-// Google Benchmark library
+#include <cuda/stream_ref>
 #include <cuda_runtime_api.h>
 
 #include <benchmark/benchmark.h>
@@ -68,7 +66,7 @@ class cuda_event_timer {
    */
   cuda_event_timer(benchmark::State& state,
                    bool flush_l2_cache,
-                   rmm::cuda_stream_view stream = rmm::cuda_stream_default);
+                   cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}});
 
   // The user will HAVE to provide a benchmark::State object to set
   // the timer so we disable the default c'tor.
@@ -88,6 +86,6 @@ class cuda_event_timer {
  private:
   cudaEvent_t start{};
   cudaEvent_t stop{};
-  rmm::cuda_stream_view stream{};
+  cuda::stream_ref stream{cudaStream_t{nullptr}};
   benchmark::State* p_state{};
 };
diff --git a/cpp/include/rmm/detail/format.hpp b/cpp/include/rmm/detail/format.hpp
index 24d8bed70..d79749c3b 100644
--- a/cpp/include/rmm/detail/format.hpp
+++ b/cpp/include/rmm/detail/format.hpp
@@ -5,7 +5,7 @@
 
 #pragma once
 
-#include <rmm/cuda_stream_view.hpp>
+#include <cuda/stream_ref>
 
 #include <array>
 #include <cstdio>
@@ -31,10 +31,10 @@ inline std::string format_bytes(std::size_t value)
 }
 
 // Stringify a stream ID
-inline std::string format_stream(rmm::cuda_stream_view stream)
+inline std::string format_stream(cuda::stream_ref stream)
 {
   std::stringstream sstr{};
-  sstr << std::hex << stream.value();
+  sstr << std::hex << stream.get();
   return sstr.str();
 }
 
diff --git a/cpp/include/rmm/device_buffer.hpp b/cpp/include/rmm/device_buffer.hpp
index 91514cd9d..c21cde222 100644
--- a/cpp/include/rmm/device_buffer.hpp
+++ b/cpp/include/rmm/device_buffer.hpp
@@ -13,6 +13,7 @@
 #include <rmm/resource_ref.hpp>
 
 #include <cuda/memory_resource>
+#include <cuda/stream_ref>
 #include <cuda_runtime_api.h>
 
 #include <cassert>
@@ -379,8 +380,9 @@ class device_buffer {
   void* _data{nullptr};  ///< Pointer to device memory allocation
   std::size_t _size{};   ///< Requested size of the device memory allocation
   std::size_t _alignment{rmm::CUDA_ALLOCATION_ALIGNMENT};  ///< The alignment of the allocation
-  std::size_t _capacity{};     ///< The actual size of the device memory allocation
-  cuda_stream_view _stream{};  ///< Stream to use for device memory deallocation
+  std::size_t _capacity{};  ///< The actual size of the device memory allocation
+  cuda::stream_ref _stream{
+    cudaStream_t{nullptr}};  ///< Stream to use for device memory deallocation
 
   cuda::mr::any_resource<cuda::mr::device_accessible> _mr;  ///< The memory resource used to
                                                             ///< allocate/deallocate device memory
diff --git a/cpp/include/rmm/device_scalar.hpp b/cpp/include/rmm/device_scalar.hpp
index 45e58cb51..da6949264 100644
--- a/cpp/include/rmm/device_scalar.hpp
+++ b/cpp/include/rmm/device_scalar.hpp
@@ -11,6 +11,8 @@
 #include <rmm/mr/per_device_resource.hpp>
 #include <rmm/resource_ref.hpp>
 
+#include <cuda/stream_ref>
+
 #include <type_traits>
 
 namespace RMM_NAMESPACE {
@@ -84,7 +86,7 @@ class device_scalar {
    * @param mr Optional, resource with which to allocate.
    */
   explicit device_scalar(
-    cuda_stream_view stream,
+    cuda::stream_ref stream,
     cuda::mr::any_resource<cuda::mr::device_accessible> mr = mr::get_current_device_resource_ref())
     : _storage{1, stream, std::move(mr)}
   {
@@ -110,7 +112,7 @@ class device_scalar {
    */
   explicit device_scalar(
     value_type const& initial_value,
-    cuda_stream_view stream,
+    cuda::stream_ref stream,
     cuda::mr::any_resource<cuda::mr::device_accessible> mr = mr::get_current_device_resource_ref())
     : _storage{1, stream, std::move(mr)}
   {
@@ -131,7 +133,7 @@ class device_scalar {
    */
   device_scalar(
     device_scalar const& other,
-    cuda_stream_view stream,
+    cuda::stream_ref stream,
     cuda::mr::any_resource<cuda::mr::device_accessible> mr = mr::get_current_device_resource_ref())
     : _storage{other._storage, stream, std::move(mr)}
   {
@@ -153,7 +155,7 @@ class device_scalar {
    * @return T The value of the scalar.
    * @param stream CUDA stream on which to perform the copy and synchronize.
    */
-  [[nodiscard]] value_type value(cuda_stream_view stream) const
+  [[nodiscard]] value_type value(cuda::stream_ref stream) const
   {
     return _storage.front_element(stream);
   }
@@ -191,14 +193,14 @@ class device_scalar {
    * @param value The host value which will be copied to device
    * @param stream CUDA stream on which to perform the copy
    */
-  void set_value_async(value_type const& value, cuda_stream_view stream)
+  void set_value_async(value_type const& value, cuda::stream_ref stream)
   {
     _storage.set_element_async(0, value, stream);
   }
 
   // Disallow passing literals to set_value to avoid race conditions where the memory holding the
   // literal can be freed before the async memcpy / memset executes.
-  void set_value_async(value_type&&, cuda_stream_view) = delete;
+  void set_value_async(value_type&&, cuda::stream_ref) = delete;
 
   /**
    * @brief Sets the value of the `device_scalar` to zero on the specified stream.
@@ -214,7 +216,7 @@ class device_scalar {
    *
    * @param stream CUDA stream on which to perform the copy
    */
-  void set_value_to_zero_async(cuda_stream_view stream)
+  void set_value_to_zero_async(cuda::stream_ref stream)
   {
     _storage.set_element_to_zero_async(value_type{0}, stream);
   }
@@ -261,7 +263,7 @@ class device_scalar {
    *
    * @param stream Stream to be used for deallocation
    */
-  void set_stream(cuda_stream_view stream) noexcept { _storage.set_stream(stream); }
+  void set_stream(cuda::stream_ref stream) noexcept { _storage.set_stream(stream); }
 
  private:
   rmm::device_uvector<T> _storage;
diff --git a/cpp/include/rmm/device_uvector.hpp b/cpp/include/rmm/device_uvector.hpp
index fb2651bbf..52bc1cda7 100644
--- a/cpp/include/rmm/device_uvector.hpp
+++ b/cpp/include/rmm/device_uvector.hpp
@@ -15,6 +15,7 @@
 
 #include <cuda/std/iterator>
 #include <cuda/std/span>
+#include <cuda/stream_ref>
 
 #include <cstddef>
 #include <type_traits>
@@ -126,7 +127,7 @@ class device_uvector {
    */
   explicit device_uvector(
     size_type size,
-    cuda_stream_view stream,
+    cuda::stream_ref stream,
     cuda::mr::any_resource<cuda::mr::device_accessible> mr = mr::get_current_device_resource_ref())
     : _storage{elements_to_bytes(size), std::alignment_of_v<T>, stream, std::move(mr)}
   {
@@ -143,7 +144,7 @@ class device_uvector {
    */
   explicit device_uvector(
     device_uvector const& other,
-    cuda_stream_view stream,
+    cuda::stream_ref stream,
     cuda::mr::any_resource<cuda::mr::device_accessible> mr = mr::get_current_device_resource_ref())
     : _storage{other._storage, stream, std::move(mr)}
   {
@@ -210,17 +211,17 @@ class device_uvector {
    * @param value The value to copy to the specified element
    * @param stream The stream on which to perform the copy
    */
-  void set_element_async(size_type element_index, value_type const& value, cuda_stream_view stream)
+  void set_element_async(size_type element_index, value_type const& value, cuda::stream_ref stream)
   {
     RMM_EXPECTS(
       element_index < size(), "Attempt to access out of bounds element.", rmm::out_of_range);
     RMM_CUDA_TRY(cudaMemcpyAsync(
-      element_ptr(element_index), &value, sizeof(value), cudaMemcpyDefault, stream.value()));
+      element_ptr(element_index), &value, sizeof(value), cudaMemcpyDefault, stream.get()));
   }
 
   // We delete the r-value reference overload to prevent asynchronously copying from a literal or
   // implicit temporary value after it is deleted or goes out of scope.
-  void set_element_async(size_type, value_type const&&, cuda_stream_view) = delete;
+  void set_element_async(size_type, value_type const&&, cuda::stream_ref) = delete;
 
   /**
    * @brief Asynchronously sets the specified element to zero in device memory.
@@ -244,12 +245,11 @@ class device_uvector {
    * @param element_index Index of the target element
    * @param stream The stream on which to perform the copy
    */
-  void set_element_to_zero_async(size_type element_index, cuda_stream_view stream)
+  void set_element_to_zero_async(size_type element_index, cuda::stream_ref stream)
   {
     RMM_EXPECTS(
       element_index < size(), "Attempt to access out of bounds element.", rmm::out_of_range);
-    RMM_CUDA_TRY(
-      cudaMemsetAsync(element_ptr(element_index), 0, sizeof(value_type), stream.value()));
+    RMM_CUDA_TRY(cudaMemsetAsync(element_ptr(element_index), 0, sizeof(value_type), stream.get()));
   }
 
   /**
@@ -281,10 +281,10 @@ class device_uvector {
    * @param value The value to copy to the specified element
    * @param stream The stream on which to perform the copy
    */
-  void set_element(size_type element_index, T const& value, cuda_stream_view stream)
+  void set_element(size_type element_index, T const& value, cuda::stream_ref stream)
   {
     set_element_async(element_index, value, stream);
-    stream.synchronize_no_throw();
+    RMM_ASSERT_CUDA_SUCCESS(cudaStreamSynchronize(stream.get()));
   }
 
   /**
@@ -299,14 +299,14 @@ class device_uvector {
    * @param stream The stream on which to perform the copy
    * @return The value of the specified element
    */
-  [[nodiscard]] value_type element(size_type element_index, cuda_stream_view stream) const
+  [[nodiscard]] value_type element(size_type element_index, cuda::stream_ref stream) const
   {
     RMM_EXPECTS(
       element_index < size(), "Attempt to access out of bounds element.", rmm::out_of_range);
     value_type value;
     RMM_CUDA_TRY(cudaMemcpyAsync(
-      &value, element_ptr(element_index), sizeof(value), cudaMemcpyDefault, stream.value()));
-    stream.synchronize();
+      &value, element_ptr(element_index), sizeof(value), cudaMemcpyDefault, stream.get()));
+    RMM_CUDA_TRY(cudaStreamSynchronize(stream.get()));
     return value;
   }
 
@@ -321,7 +321,7 @@ class device_uvector {
    * @param stream The stream on which to perform the copy
    * @return The value of the first element
    */
-  [[nodiscard]] value_type front_element(cuda_stream_view stream) const
+  [[nodiscard]] value_type front_element(cuda::stream_ref stream) const
   {
     return element(0, stream);
   }
@@ -337,7 +337,7 @@ class device_uvector {
    * @param stream The stream on which to perform the copy
    * @return The value of the last element
    */
-  [[nodiscard]] value_type back_element(cuda_stream_view stream) const
+  [[nodiscard]] value_type back_element(cuda::stream_ref stream) const
   {
     return element(size() - 1, stream);
   }
@@ -354,7 +354,7 @@ class device_uvector {
    * @param new_capacity The desired capacity (number of elements)
    * @param stream The stream on which to perform the allocation/copy (if any)
    */
-  void reserve(size_type new_capacity, cuda_stream_view stream)
+  void reserve(size_type new_capacity, cuda::stream_ref stream)
   {
     _storage.reserve(elements_to_bytes(new_capacity), stream);
   }
@@ -375,7 +375,7 @@ class device_uvector {
    * @param new_size The desired number of elements
    * @param stream The stream on which to perform the allocation/copy (if any)
    */
-  void resize(size_type new_size, cuda_stream_view stream)
+  void resize(size_type new_size, cuda::stream_ref stream)
   {
     _storage.resize(elements_to_bytes(new_size), stream);
   }
@@ -387,7 +387,7 @@ class device_uvector {
    *
    * @param stream Stream on which to perform allocation and copy
    */
-  void shrink_to_fit(cuda_stream_view stream) { _storage.shrink_to_fit(stream); }
+  void shrink_to_fit(cuda::stream_ref stream) { _storage.shrink_to_fit(stream); }
 
   /**
    * @brief Release ownership of device memory storage.
@@ -612,7 +612,7 @@ class device_uvector {
    *
    * @param stream The stream to use for deallocation
    */
-  void set_stream(cuda_stream_view stream) noexcept { _storage.set_stream(stream); }
+  void set_stream(cuda::stream_ref stream) noexcept { _storage.set_stream(stream); }
 
  private:
   device_buffer _storage{};  ///< Device memory storage for vector elements
diff --git a/cpp/include/rmm/mr/detail/arena.hpp b/cpp/include/rmm/mr/detail/arena.hpp
index ec5f501b3..99e6be372 100644
--- a/cpp/include/rmm/mr/detail/arena.hpp
+++ b/cpp/include/rmm/mr/detail/arena.hpp
@@ -7,7 +7,6 @@
 
 #include <rmm/aligned.hpp>
 #include <rmm/cuda_device.hpp>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/detail/error.hpp>
 #include <rmm/detail/export.hpp>
 #include <rmm/detail/format.hpp>
@@ -16,6 +15,7 @@
 #include <rmm/resource_ref.hpp>
 
 #include <cuda/memory_resource>
+#include <cuda/stream_ref>
 #include <cuda_runtime_api.h>
 
 #include <algorithm>
@@ -595,10 +595,10 @@ class global_arena final {
    * that was passed to the `allocate` call that returned `ptr`.
    * @return bool true if the allocation is found, false otherwise.
    */
-  bool deallocate(cuda_stream_view stream, void* ptr, std::size_t size)
+  bool deallocate(cuda::stream_ref stream, void* ptr, std::size_t size)
   {
     RMM_LOGGING_ASSERT(handles(size));
-    stream.synchronize_no_throw();
+    RMM_ASSERT_CUDA_SUCCESS(cudaStreamSynchronize(stream.get()));
     return deallocate_sync(ptr, size);
   }
 
@@ -818,7 +818,7 @@ class arena {
    * that was passed to the `allocate` call that returned `ptr`.
    * @return bool true if the allocation is found, false otherwise.
    */
-  bool deallocate(cuda_stream_view stream, void* ptr, std::size_t size)
+  bool deallocate(cuda::stream_ref stream, void* ptr, std::size_t size)
   {
     if (global_arena::handles(size) && global_arena_.deallocate(stream, ptr, size)) { return true; }
     return deallocate_sync(ptr, size);
diff --git a/cpp/include/rmm/mr/detail/arena_memory_resource_impl.hpp b/cpp/include/rmm/mr/detail/arena_memory_resource_impl.hpp
index ca3175eb7..253d13b7f 100644
--- a/cpp/include/rmm/mr/detail/arena_memory_resource_impl.hpp
+++ b/cpp/include/rmm/mr/detail/arena_memory_resource_impl.hpp
@@ -4,12 +4,12 @@
  */
 #pragma once
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/detail/export.hpp>
 #include <rmm/mr/detail/arena.hpp>
 #include <rmm/resource_ref.hpp>
 
 #include <cuda/memory_resource>
+#include <cuda/stream_ref>
 
 #include <cstddef>
 #include <map>
@@ -79,15 +79,15 @@ class arena_memory_resource_impl {
 
   void defragment();
 
-  void deallocate_from_other_arena(cuda_stream_view stream, void* ptr, std::size_t bytes);
+  void deallocate_from_other_arena(cuda::stream_ref stream, void* ptr, std::size_t bytes);
 
-  arena& get_arena(cuda_stream_view stream);
+  arena& get_arena(cuda::stream_ref stream);
   arena& get_thread_arena();
-  arena& get_stream_arena(cuda_stream_view stream);
+  arena& get_stream_arena(cuda::stream_ref stream);
 
   void dump_memory_log(std::size_t bytes);
 
-  static bool use_per_thread_arena(cuda_stream_view stream);
+  static bool use_per_thread_arena(cuda::stream_ref stream);
 
   global_arena global_arena_;
   std::map<std::thread::id, std::shared_ptr<arena>> thread_arenas_;
diff --git a/cpp/include/rmm/mr/detail/failure_callback_resource_adaptor_impl.hpp b/cpp/include/rmm/mr/detail/failure_callback_resource_adaptor_impl.hpp
index 9b3e61f14..13c4c4ab2 100644
--- a/cpp/include/rmm/mr/detail/failure_callback_resource_adaptor_impl.hpp
+++ b/cpp/include/rmm/mr/detail/failure_callback_resource_adaptor_impl.hpp
@@ -5,12 +5,13 @@
 #pragma once
 
 #include <rmm/aligned.hpp>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/detail/export.hpp>
 #include <rmm/mr/failure_callback_t.hpp>
 #include <rmm/resource_ref.hpp>
 
 #include <cuda/memory_resource>
+#include <cuda/stream_ref>
+#include <cuda_runtime_api.h>
 
 #include <cstddef>
 #include <utility>
@@ -86,14 +87,14 @@ class failure_callback_resource_adaptor_impl {
 
   void* allocate_sync(std::size_t bytes, std::size_t alignment = alignof(std::max_align_t))
   {
-    return allocate(cuda_stream_view{}, bytes, alignment);
+    return allocate(cuda::stream_ref{cudaStream_t{nullptr}}, bytes, alignment);
   }
 
   void deallocate_sync(void* ptr,
                        std::size_t bytes,
                        std::size_t alignment = alignof(std::max_align_t)) noexcept
   {
-    deallocate(cuda_stream_view{}, ptr, bytes, alignment);
+    deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
   }
 
   RMM_CONSTEXPR_FRIEND void get_property(failure_callback_resource_adaptor_impl const&,
diff --git a/cpp/include/rmm/mr/detail/fixed_size_memory_resource_impl.hpp b/cpp/include/rmm/mr/detail/fixed_size_memory_resource_impl.hpp
index c651f5077..e218cfe74 100644
--- a/cpp/include/rmm/mr/detail/fixed_size_memory_resource_impl.hpp
+++ b/cpp/include/rmm/mr/detail/fixed_size_memory_resource_impl.hpp
@@ -4,12 +4,12 @@
  */
 #pragma once
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/mr/detail/fixed_size_free_list.hpp>
 #include <rmm/mr/detail/stream_ordered_memory_resource.hpp>
 #include <rmm/resource_ref.hpp>
 
 #include <cuda/memory_resource>
+#include <cuda/stream_ref>
 
 #include <cstddef>
 #include <mutex>
@@ -70,7 +70,7 @@ class fixed_size_memory_resource_impl final
 
   [[nodiscard]] std::size_t get_maximum_allocation_size() const;
 
-  block_type expand_pool(std::size_t size, free_list& blocks, cuda_stream_view stream);
+  block_type expand_pool(std::size_t size, free_list& blocks, cuda::stream_ref stream);
 
   split_block allocate_from_block(block_type const& block, std::size_t size);
 
@@ -79,7 +79,7 @@ class fixed_size_memory_resource_impl final
   std::pair<std::size_t, std::size_t> free_list_summary(free_list const& blocks);
 
  private:
-  free_list blocks_from_upstream(cuda_stream_view stream);
+  free_list blocks_from_upstream(cuda::stream_ref stream);
 
   void release();
 
diff --git a/cpp/include/rmm/mr/detail/pool_memory_resource_impl.hpp b/cpp/include/rmm/mr/detail/pool_memory_resource_impl.hpp
index 43e109064..a82afe258 100644
--- a/cpp/include/rmm/mr/detail/pool_memory_resource_impl.hpp
+++ b/cpp/include/rmm/mr/detail/pool_memory_resource_impl.hpp
@@ -4,12 +4,12 @@
  */
 #pragma once
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/mr/detail/coalescing_free_list.hpp>
 #include <rmm/mr/detail/stream_ordered_memory_resource.hpp>
 #include <rmm/resource_ref.hpp>
 
 #include <cuda/memory_resource>
+#include <cuda/stream_ref>
 
 #include <cstddef>
 #include <memory>
@@ -67,11 +67,11 @@ class pool_memory_resource_impl final
   using lock_guard = std::lock_guard<std::mutex>;
 
   [[nodiscard]] std::size_t get_maximum_allocation_size() const;
-  block_type try_to_expand(std::size_t try_size, std::size_t min_size, cuda_stream_view stream);
+  block_type try_to_expand(std::size_t try_size, std::size_t min_size, cuda::stream_ref stream);
   void initialize_pool(std::size_t initial_size, std::optional<std::size_t> maximum_size);
-  block_type expand_pool(std::size_t size, free_list& blocks, cuda_stream_view stream);
+  block_type expand_pool(std::size_t size, free_list& blocks, cuda::stream_ref stream);
   [[nodiscard]] std::size_t size_to_grow(std::size_t size) const;
-  block_type block_from_upstream(std::size_t size, cuda_stream_view stream);
+  block_type block_from_upstream(std::size_t size, cuda::stream_ref stream);
   split_block allocate_from_block(block_type const& block, std::size_t size);
   block_type free_block(void* ptr, std::size_t size) noexcept;
   void release();
diff --git a/cpp/include/rmm/mr/detail/stream_ordered_memory_resource.hpp b/cpp/include/rmm/mr/detail/stream_ordered_memory_resource.hpp
index e8de8b65e..3aa72a2b6 100644
--- a/cpp/include/rmm/mr/detail/stream_ordered_memory_resource.hpp
+++ b/cpp/include/rmm/mr/detail/stream_ordered_memory_resource.hpp
@@ -60,7 +60,7 @@ struct crtp {
  * documented separately:
  *
  * 1. `std::size_t get_maximum_allocation_size() const`
- * 2. `block_type expand_pool(std::size_t size, free_list& blocks, cuda_stream_view stream)`
+ * 2. `block_type expand_pool(std::size_t size, free_list& blocks, cuda::stream_ref stream)`
  * 3. `split_block allocate_from_block(block_type const& b, std::size_t size)`
  * 4. `block_type free_block(void* ptr, std::size_t size) noexcept`
  */
@@ -89,15 +89,13 @@ class stream_ordered_memory_resource : public crtp<PoolResource> {
    */
   void* allocate(cuda::stream_ref stream, std::size_t bytes, std::size_t /*alignment*/)
   {
-    auto const strm = cuda_stream_view{stream};
-
-    RMM_LOG_TRACE("[A][stream %s][%zuB]", rmm::detail::format_stream(strm), bytes);
+    RMM_LOG_TRACE("[A][stream %s][%zuB]", rmm::detail::format_stream(stream), bytes);
 
     if (bytes == 0) { return nullptr; }
 
     lock_guard lock(mtx_);
 
-    auto stream_event = get_event(strm);
+    auto stream_event = get_event(stream);
 
     bytes = rmm::align_up(bytes, rmm::CUDA_ALLOCATION_ALIGNMENT);
     RMM_EXPECTS(bytes <= this->underlying().get_maximum_allocation_size(),
@@ -129,14 +127,12 @@ class stream_ordered_memory_resource : public crtp<PoolResource> {
                   std::size_t bytes,
                   std::size_t /*alignment*/) noexcept
   {
-    auto const strm = cuda_stream_view{stream};
-
-    RMM_LOG_TRACE("[D][stream %s][%zuB][%p]", rmm::detail::format_stream(strm), bytes, ptr);
+    RMM_LOG_TRACE("[D][stream %s][%zuB][%p]", rmm::detail::format_stream(stream), bytes, ptr);
 
     if (bytes == 0 || ptr == nullptr) { return; }
 
     lock_guard lock(mtx_);
-    auto stream_event = get_event(strm);
+    auto stream_event = get_event(stream);
 
     bytes            = rmm::align_up(bytes, rmm::CUDA_ALLOCATION_ALIGNMENT);
     auto const block = this->underlying().free_block(ptr, bytes);
@@ -144,7 +140,7 @@ class stream_ordered_memory_resource : public crtp<PoolResource> {
     // TODO: cudaEventRecord has significant overhead on deallocations. For the non-PTDS case
     // we may be able to delay recording the event in some situations. But using events rather
     // than streams allows stealing from deleted streams.
-    RMM_ASSERT_CUDA_SUCCESS(cudaEventRecord(stream_event.event, strm.value()));
+    RMM_ASSERT_CUDA_SUCCESS(cudaEventRecord(stream_event.event, stream.get()));
 
     stream_free_blocks_[stream_event].insert(block);
 
@@ -163,9 +159,9 @@ class stream_ordered_memory_resource : public crtp<PoolResource> {
   [[nodiscard]] void* allocate_sync(std::size_t bytes,
                                     std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT)
   {
-    auto const stream = cuda_stream_view{};
+    auto const stream = cuda::stream_ref{cudaStream_t{nullptr}};
     void* ptr         = allocate(stream, bytes, alignment);
-    stream.synchronize();
+    RMM_CUDA_TRY(cudaStreamSynchronize(stream.get()));
     return ptr;
   }
 
@@ -181,7 +177,7 @@ class stream_ordered_memory_resource : public crtp<PoolResource> {
     std::size_t bytes,
     [[maybe_unused]] std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept
   {
-    deallocate(cuda_stream_view{}, ptr, bytes, alignment);
+    deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
   }
 
  protected:
@@ -217,7 +213,7 @@ class stream_ordered_memory_resource : public crtp<PoolResource> {
    * @param stream The stream on which the memory is to be used.
    * @return block_type a block of at least `size` bytes
    */
-  // block_type expand_pool(std::size_t size, free_list& blocks, cuda_stream_view stream)
+  // block_type expand_pool(std::size_t size, free_list& blocks, cuda::stream_ref stream)
 
   /// Pair representing a block that has been split for allocation
   using split_block = std::pair<block_type, block_type>;
@@ -252,12 +248,12 @@ class stream_ordered_memory_resource : public crtp<PoolResource> {
    * @param block The block to insert into the pool.
    * @param stream The stream on which the memory was last used.
    */
-  void insert_block(block_type const& block, cuda_stream_view stream)
+  void insert_block(block_type const& block, cuda::stream_ref stream)
   {
     stream_free_blocks_[get_event(stream)].insert(block);
   }
 
-  void insert_blocks(free_list&& blocks, cuda_stream_view stream)
+  void insert_blocks(free_list&& blocks, cuda::stream_ref stream)
   {
     stream_free_blocks_[get_event(stream)].insert(std::move(blocks));
   }
@@ -301,9 +297,27 @@ class stream_ordered_memory_resource : public crtp<PoolResource> {
    * @param stream The stream for which to get an event.
    * @return The stream_event for `stream`.
    */
-  stream_event_pair get_event(cuda_stream_view stream)
+  static bool is_per_thread_default(cuda::stream_ref stream) noexcept
+  {
+#ifdef CUDA_API_PER_THREAD_DEFAULT_STREAM
+    return stream.get() == cudaStreamPerThread || stream.get() == nullptr;
+#else
+    return stream.get() == cudaStreamPerThread;
+#endif
+  }
+
+  static bool is_default(cuda::stream_ref stream) noexcept
+  {
+#ifdef CUDA_API_PER_THREAD_DEFAULT_STREAM
+    return stream.get() == cudaStreamLegacy;
+#else
+    return stream.get() == cudaStreamLegacy || stream.get() == nullptr;
+#endif
+  }
+
+  stream_event_pair get_event(cuda::stream_ref stream)
   {
-    if (stream.is_per_thread_default()) {
+    if (is_per_thread_default(stream)) {
       // Create a thread-local event for each device. These events are
       // deliberately leaked since the destructor needs to call into
       // the CUDA runtime and thread_local destructors (can) run below
@@ -320,7 +334,7 @@ class stream_ordered_memory_resource : public crtp<PoolResource> {
         }
         return e;
       }();
-      return stream_event_pair{stream.value(), event};
+      return stream_event_pair{stream.get(), event};
     }
     // We use cudaStreamLegacy as the event map key for the default stream for consistency between
     // PTDS and non-PTDS mode. In PTDS mode, the cudaStreamLegacy map key will only exist if the
@@ -328,7 +342,7 @@ class stream_ordered_memory_resource : public crtp<PoolResource> {
     // at construction. For consistency, the same key is used for null stream free lists in
     // non-PTDS mode.
     // NOLINTNEXTLINE(cppcoreguidelines-pro-type-cstyle-cast)
-    auto* const stream_to_store = stream.is_default() ? cudaStreamLegacy : stream.value();
+    auto* const stream_to_store = is_default(stream) ? cudaStreamLegacy : stream.get();
     stream_id_type stream_id{};
     RMM_ASSERT_CUDA_SUCCESS(cudaStreamGetId(stream_to_store, &stream_id));
     auto const iter = stream_events_.find(stream_id);
@@ -393,7 +407,7 @@ class stream_ordered_memory_resource : public crtp<PoolResource> {
 
     // no large enough blocks available after merging, so grow the pool
     block_type const block =
-      this->underlying().expand_pool(size, blocks, cuda_stream_view{stream_event.stream});
+      this->underlying().expand_pool(size, blocks, cuda::stream_ref{stream_event.stream});
 
     return allocate_and_insert_remainder(block, size, blocks);
   }
diff --git a/cpp/include/rmm/mr/polymorphic_allocator.hpp b/cpp/include/rmm/mr/polymorphic_allocator.hpp
index 4e2559deb..0d47db73e 100644
--- a/cpp/include/rmm/mr/polymorphic_allocator.hpp
+++ b/cpp/include/rmm/mr/polymorphic_allocator.hpp
@@ -12,6 +12,7 @@
 #include <rmm/resource_ref.hpp>
 
 #include <cuda/memory_resource>
+#include <cuda/stream_ref>
 
 #include <cstddef>
 #include <memory>
@@ -251,8 +252,9 @@ class stream_allocator_adaptor {
   [[nodiscard]] Allocator underlying_allocator() const noexcept { return alloc_; }
 
  private:
-  Allocator alloc_;          ///< Underlying allocator used for (de)allocation
-  cuda_stream_view stream_;  ///< Stream on which (de)allocations are performed
+  Allocator alloc_;  ///< Underlying allocator used for (de)allocation
+  cuda::stream_ref stream_{
+    cudaStream_t{nullptr}};  ///< Stream on which (de)allocations are performed
 };
 
 /**
diff --git a/cpp/include/rmm/mr/thrust_allocator_adaptor.hpp b/cpp/include/rmm/mr/thrust_allocator_adaptor.hpp
index bf4d5024f..c435c3471 100644
--- a/cpp/include/rmm/mr/thrust_allocator_adaptor.hpp
+++ b/cpp/include/rmm/mr/thrust_allocator_adaptor.hpp
@@ -14,6 +14,7 @@
 #include <rmm/resource_ref.hpp>
 
 #include <cuda/memory_resource>
+#include <cuda/stream_ref>
 #include <thrust/device_malloc_allocator.h>
 #include <thrust/device_ptr.h>
 #include <thrust/memory.h>
@@ -179,7 +180,7 @@ class thrust_allocator : public thrust::device_malloc_allocator<T> {
   }
 
  private:
-  cuda_stream_view _stream{};
+  cuda::stream_ref _stream{cudaStream_t{nullptr}};
   mutable cuda::mr::any_resource<cuda::mr::device_accessible> _mr{
     rmm::mr::get_current_device_resource_ref()};
   cuda_device_id _device{get_current_cuda_device()};
diff --git a/cpp/include/rmm/prefetch.hpp b/cpp/include/rmm/prefetch.hpp
index 2782cd144..1ebc5f96a 100644
--- a/cpp/include/rmm/prefetch.hpp
+++ b/cpp/include/rmm/prefetch.hpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -12,6 +12,7 @@
 #include <rmm/error.hpp>
 
 #include <cuda/std/span>
+#include <cuda/stream_ref>
 
 namespace RMM_NAMESPACE {
 
diff --git a/cpp/src/device_buffer.cpp b/cpp/src/device_buffer.cpp
index 102f5c701..ed1c2bfb7 100644
--- a/cpp/src/device_buffer.cpp
+++ b/cpp/src/device_buffer.cpp
@@ -8,6 +8,7 @@
 #include <rmm/device_buffer.hpp>
 #include <rmm/error.hpp>
 
+#include <cuda/stream_ref>
 #include <cuda_runtime_api.h>
 
 #include <memory>
@@ -80,8 +81,8 @@ device_buffer::device_buffer(device_buffer&& other) noexcept
   other._size      = 0;
   other._alignment = 1;
   other._capacity  = 0;
-  other.set_stream(cuda_stream_view{});
-  other._device = cuda_device_id{-1};
+  other._stream    = cuda::stream_ref{cudaStream_t{nullptr}};
+  other._device    = cuda_device_id{-1};
 }
 
 device_buffer& device_buffer::operator=(device_buffer&& other) noexcept
@@ -102,8 +103,8 @@ device_buffer& device_buffer::operator=(device_buffer&& other) noexcept
     other._size      = 0;
     other._alignment = 1;
     other._capacity  = 0;
-    other.set_stream(cuda_stream_view{});
-    other._device = cuda_device_id{-1};
+    other._stream    = cuda::stream_ref{cudaStream_t{nullptr}};
+    other._device    = cuda_device_id{-1};
   }
   return *this;
 }
@@ -112,19 +113,19 @@ device_buffer::~device_buffer() noexcept
 {
   cuda_set_device_raii dev{_device};
   deallocate_async();
-  _stream = cuda_stream_view{};
+  _stream = cuda::stream_ref{cudaStream_t{nullptr}};
 }
 
 void device_buffer::allocate_async(std::size_t bytes)
 {
   _size     = bytes;
   _capacity = bytes;
-  _data     = (bytes > 0) ? _mr.allocate(stream(), bytes, alignment()) : nullptr;
+  _data     = (bytes > 0) ? _mr.allocate(_stream, bytes, alignment()) : nullptr;
 }
 
 void device_buffer::deallocate_async() noexcept
 {
-  if (capacity() > 0) { _mr.deallocate(stream(), data(), capacity(), alignment()); }
+  if (capacity() > 0) { _mr.deallocate(_stream, data(), capacity(), alignment()); }
   _size      = 0;
   _alignment = 1;
   _capacity  = 0;
@@ -137,7 +138,7 @@ void device_buffer::copy_async(void const* source, std::size_t bytes)
     RMM_EXPECTS(nullptr != source, "Invalid copy from nullptr.");
     RMM_EXPECTS(nullptr != _data, "Invalid copy to nullptr.");
 
-    RMM_CUDA_TRY(cudaMemcpyAsync(_data, source, bytes, cudaMemcpyDefault, stream().value()));
+    RMM_CUDA_TRY(cudaMemcpyAsync(_data, source, bytes, cudaMemcpyDefault, _stream.get()));
   }
 }
 
@@ -148,7 +149,7 @@ void device_buffer::reserve(std::size_t new_capacity, cuda_stream_view stream)
     cuda_set_device_raii dev{_device};
     auto tmp            = device_buffer{new_capacity, alignment(), stream, _mr};
     auto const old_size = size();
-    RMM_CUDA_TRY(cudaMemcpyAsync(tmp.data(), data(), size(), cudaMemcpyDefault, stream.value()));
+    RMM_CUDA_TRY(cudaMemcpyAsync(tmp.data(), data(), size(), cudaMemcpyDefault, _stream.get()));
     *this = std::move(tmp);
     _size = old_size;
   }
@@ -164,7 +165,7 @@ void device_buffer::resize(std::size_t new_size, cuda_stream_view stream)
   } else {
     cuda_set_device_raii dev{_device};
     auto tmp = device_buffer{new_size, alignment(), stream, _mr};
-    RMM_CUDA_TRY(cudaMemcpyAsync(tmp.data(), data(), size(), cudaMemcpyDefault, stream.value()));
+    RMM_CUDA_TRY(cudaMemcpyAsync(tmp.data(), data(), size(), cudaMemcpyDefault, _stream.get()));
     *this = std::move(tmp);
   }
 }
diff --git a/cpp/src/exec_policy.cpp b/cpp/src/exec_policy.cpp
index fe7baea59..f5b36b8d1 100644
--- a/cpp/src/exec_policy.cpp
+++ b/cpp/src/exec_policy.cpp
@@ -9,8 +9,8 @@ namespace rmm {
 
 exec_policy::exec_policy(cuda_stream_view stream,
                          cuda::mr::any_resource<cuda::mr::device_accessible> mr)
-  : thrust_exec_policy_t(
-      thrust::cuda::par(mr::thrust_allocator<char>(stream, std::move(mr))).on(stream.value()))
+  : thrust_exec_policy_t(thrust::cuda::par(mr::thrust_allocator<char>(stream, std::move(mr)))
+                           .on(cuda::stream_ref{stream}.get()))
 {
 }
 
@@ -18,7 +18,7 @@ exec_policy_nosync::exec_policy_nosync(cuda_stream_view stream,
                                        cuda::mr::any_resource<cuda::mr::device_accessible> mr)
   : thrust_exec_policy_nosync_t(
       thrust::cuda::par_nosync(mr::thrust_allocator<char>(stream, std::move(mr)))
-        .on(stream.value()))
+        .on(cuda::stream_ref{stream}.get()))
 {
 }
 
diff --git a/cpp/src/mr/detail/aligned_resource_adaptor_impl.cpp b/cpp/src/mr/detail/aligned_resource_adaptor_impl.cpp
index b99dc5750..0556ff93a 100644
--- a/cpp/src/mr/detail/aligned_resource_adaptor_impl.cpp
+++ b/cpp/src/mr/detail/aligned_resource_adaptor_impl.cpp
@@ -4,10 +4,12 @@
  */
 
 #include <rmm/aligned.hpp>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/detail/error.hpp>
 #include <rmm/mr/detail/aligned_resource_adaptor_impl.hpp>
 
+#include <cuda/stream_ref>
+#include <cuda_runtime_api.h>
+
 #include <algorithm>
 #include <cstddef>
 
@@ -89,14 +91,14 @@ void aligned_resource_adaptor_impl::deallocate(cuda::stream_ref stream,
 
 void* aligned_resource_adaptor_impl::allocate_sync(std::size_t bytes, std::size_t alignment)
 {
-  return allocate(cuda_stream_view{}, bytes, alignment);
+  return allocate(cuda::stream_ref{cudaStream_t{nullptr}}, bytes, alignment);
 }
 
 void aligned_resource_adaptor_impl::deallocate_sync(void* ptr,
                                                     std::size_t bytes,
                                                     std::size_t alignment) noexcept
 {
-  deallocate(cuda_stream_view{}, ptr, bytes, alignment);
+  deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
 }
 
 }  // namespace detail
diff --git a/cpp/src/mr/detail/arena_memory_resource_impl.cpp b/cpp/src/mr/detail/arena_memory_resource_impl.cpp
index d3f53eb8b..72106b932 100644
--- a/cpp/src/mr/detail/arena_memory_resource_impl.cpp
+++ b/cpp/src/mr/detail/arena_memory_resource_impl.cpp
@@ -4,6 +4,7 @@
  */
 
 #include <rmm/aligned.hpp>
+#include <rmm/detail/error.hpp>
 #include <rmm/detail/format.hpp>
 #include <rmm/detail/logging_assert.hpp>
 #include <rmm/logger.hpp>
@@ -33,13 +34,12 @@ void* arena_memory_resource_impl::allocate(cuda::stream_ref stream,
                                            std::size_t /*alignment*/)
 {
   if (bytes == 0) { return nullptr; }
-  cuda_stream_view sv{stream.get()};
 #ifdef RMM_ARENA_USE_SIZE_CLASSES
   bytes = rmm::mr::detail::arena::align_to_size_class(bytes);
 #else
   bytes = rmm::align_up(bytes, rmm::CUDA_ALLOCATION_ALIGNMENT);
 #endif
-  auto& arena = get_arena(sv);
+  auto& arena = get_arena(stream);
 
   {
     std::shared_lock lock(mtx_);
@@ -67,36 +67,35 @@ void arena_memory_resource_impl::deallocate(cuda::stream_ref stream,
                                             std::size_t /*alignment*/) noexcept
 {
   if (ptr == nullptr || bytes == 0) { return; }
-  cuda_stream_view sv{stream.get()};
 #ifdef RMM_ARENA_USE_SIZE_CLASSES
   bytes = rmm::mr::detail::arena::align_to_size_class(bytes);
 #else
   bytes = rmm::align_up(bytes, rmm::CUDA_ALLOCATION_ALIGNMENT);
 #endif
-  auto& arena = get_arena(sv);
+  auto& arena = get_arena(stream);
 
   {
     std::shared_lock lock(mtx_);
-    if (arena.deallocate(sv, ptr, bytes)) { return; }
+    if (arena.deallocate(stream, ptr, bytes)) { return; }
   }
 
   {
-    sv.synchronize_no_throw();
+    RMM_ASSERT_CUDA_SUCCESS(cudaStreamSynchronize(stream.get()));
     std::unique_lock lock(mtx_);
-    deallocate_from_other_arena(sv, ptr, bytes);
+    deallocate_from_other_arena(stream, ptr, bytes);
   }
 }
 
 void* arena_memory_resource_impl::allocate_sync(std::size_t bytes, std::size_t alignment)
 {
-  return allocate(cuda_stream_view{}, bytes, alignment);
+  return allocate(cuda::stream_ref{cudaStream_t{nullptr}}, bytes, alignment);
 }
 
 void arena_memory_resource_impl::deallocate_sync(void* ptr,
                                                  std::size_t bytes,
                                                  std::size_t alignment) noexcept
 {
-  deallocate(cuda_stream_view{}, ptr, bytes, alignment);
+  deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
 }
 
 void arena_memory_resource_impl::defragment()
@@ -110,7 +109,7 @@ void arena_memory_resource_impl::defragment()
   }
 }
 
-void arena_memory_resource_impl::deallocate_from_other_arena(cuda_stream_view stream,
+void arena_memory_resource_impl::deallocate_from_other_arena(cuda::stream_ref stream,
                                                              void* ptr,
                                                              std::size_t bytes)
 {
@@ -138,7 +137,7 @@ void arena_memory_resource_impl::deallocate_from_other_arena(cuda_stream_view st
   }
 }
 
-arena_memory_resource_impl::arena& arena_memory_resource_impl::get_arena(cuda_stream_view stream)
+arena_memory_resource_impl::arena& arena_memory_resource_impl::get_arena(cuda::stream_ref stream)
 {
   if (use_per_thread_arena(stream)) { return get_thread_arena(); }
   return get_stream_arena(stream);
@@ -162,18 +161,18 @@ arena_memory_resource_impl::arena& arena_memory_resource_impl::get_thread_arena(
 }
 
 arena_memory_resource_impl::arena& arena_memory_resource_impl::get_stream_arena(
-  cuda_stream_view stream)
+  cuda::stream_ref stream)
 {
   RMM_LOGGING_ASSERT(!use_per_thread_arena(stream));
   {
     std::shared_lock lock(map_mtx_);
-    auto const iter = stream_arenas_.find(stream.value());
+    auto const iter = stream_arenas_.find(stream.get());
     if (iter != stream_arenas_.end()) { return iter->second; }
   }
   {
     std::unique_lock lock(map_mtx_);
-    stream_arenas_.emplace(stream.value(), global_arena_);
-    return stream_arenas_.at(stream.value());
+    stream_arenas_.emplace(stream.get(), global_arena_);
+    return stream_arenas_.at(stream.get());
   }
 }
 
@@ -187,9 +186,13 @@ void arena_memory_resource_impl::dump_memory_log(std::size_t bytes)
   logger_->flush();
 }
 
-bool arena_memory_resource_impl::use_per_thread_arena(cuda_stream_view stream)
+bool arena_memory_resource_impl::use_per_thread_arena(cuda::stream_ref stream)
 {
-  return stream.is_per_thread_default();
+#ifdef CUDA_API_PER_THREAD_DEFAULT_STREAM
+  return stream.get() == cudaStreamPerThread || stream.get() == nullptr;
+#else
+  return stream.get() == cudaStreamPerThread;
+#endif
 }
 
 }  // namespace detail
diff --git a/cpp/src/mr/detail/binning_memory_resource_impl.cpp b/cpp/src/mr/detail/binning_memory_resource_impl.cpp
index ff91937db..5d6dbbebd 100644
--- a/cpp/src/mr/detail/binning_memory_resource_impl.cpp
+++ b/cpp/src/mr/detail/binning_memory_resource_impl.cpp
@@ -80,14 +80,16 @@ void binning_memory_resource_impl::deallocate(cuda::stream_ref stream,
 void* binning_memory_resource_impl::allocate_sync(std::size_t bytes, std::size_t alignment)
 {
   if (bytes == 0) { return nullptr; }
-  return get_resource_ref(bytes).allocate(cuda_stream_view{}, bytes, alignment);
+  return get_resource_ref(bytes).allocate(
+    cuda::stream_ref{cudaStream_t{nullptr}}, bytes, alignment);
 }
 
 void binning_memory_resource_impl::deallocate_sync(void* ptr,
                                                    std::size_t bytes,
                                                    std::size_t alignment) noexcept
 {
-  get_resource_ref(bytes).deallocate(cuda_stream_view{}, ptr, bytes, alignment);
+  get_resource_ref(bytes).deallocate(
+    cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
 }
 
 }  // namespace detail
diff --git a/cpp/src/mr/detail/callback_memory_resource_impl.cpp b/cpp/src/mr/detail/callback_memory_resource_impl.cpp
index 2f82ff7bf..b5b42b5e0 100644
--- a/cpp/src/mr/detail/callback_memory_resource_impl.cpp
+++ b/cpp/src/mr/detail/callback_memory_resource_impl.cpp
@@ -5,6 +5,9 @@
 
 #include <rmm/mr/detail/callback_memory_resource_impl.hpp>
 
+#include <cuda/stream_ref>
+#include <cuda_runtime_api.h>
+
 #include <utility>
 
 namespace RMM_NAMESPACE {
@@ -40,14 +43,14 @@ void callback_memory_resource_impl::deallocate(cuda::stream_ref stream,
 
 void* callback_memory_resource_impl::allocate_sync(std::size_t bytes, std::size_t alignment)
 {
-  return allocate(cuda_stream_view{}, bytes, alignment);
+  return allocate(cuda::stream_ref{cudaStream_t{nullptr}}, bytes, alignment);
 }
 
 void callback_memory_resource_impl::deallocate_sync(void* ptr,
                                                     std::size_t bytes,
                                                     std::size_t alignment) noexcept
 {
-  deallocate(cuda_stream_view{}, ptr, bytes, alignment);
+  deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
 }
 
 }  // namespace detail
diff --git a/cpp/src/mr/detail/fixed_size_memory_resource_impl.cpp b/cpp/src/mr/detail/fixed_size_memory_resource_impl.cpp
index e17f6ecdd..58b3fad6c 100644
--- a/cpp/src/mr/detail/fixed_size_memory_resource_impl.cpp
+++ b/cpp/src/mr/detail/fixed_size_memory_resource_impl.cpp
@@ -4,11 +4,12 @@
  */
 
 #include <rmm/aligned.hpp>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/detail/logging_assert.hpp>
 #include <rmm/mr/detail/fixed_size_memory_resource_impl.hpp>
 
 #include <cuda/iterator>
+#include <cuda/stream_ref>
+#include <cuda_runtime_api.h>
 
 #include <cstddef>
 #include <mutex>
@@ -33,7 +34,8 @@ fixed_size_memory_resource_impl::fixed_size_memory_resource_impl(
     block_size_{align_up(block_size, rmm::CUDA_ALLOCATION_ALIGNMENT)},
     upstream_chunk_size_{block_size_ * blocks_to_preallocate}
 {
-  this->insert_blocks(std::move(blocks_from_upstream(cuda_stream_legacy)), cuda_stream_legacy);
+  this->insert_blocks(std::move(blocks_from_upstream(cuda::stream_ref{cudaStreamLegacy})),
+                      cuda::stream_ref{cudaStreamLegacy});
 }
 
 fixed_size_memory_resource_impl::~fixed_size_memory_resource_impl() { release(); }
@@ -52,14 +54,14 @@ std::size_t fixed_size_memory_resource_impl::get_maximum_allocation_size() const
 }
 
 fixed_size_memory_resource_impl::block_type fixed_size_memory_resource_impl::expand_pool(
-  std::size_t size, free_list& blocks, cuda_stream_view stream)
+  std::size_t size, free_list& blocks, cuda::stream_ref stream)
 {
   blocks.insert(std::move(blocks_from_upstream(stream)));
   return blocks.get_block(size);
 }
 
 fixed_size_memory_resource_impl::free_list fixed_size_memory_resource_impl::blocks_from_upstream(
-  cuda_stream_view stream)
+  cuda::stream_ref stream)
 {
   void* ptr = upstream_mr_.allocate(stream, upstream_chunk_size_, rmm::CUDA_ALLOCATION_ALIGNMENT);
   block_type block{ptr};
diff --git a/cpp/src/mr/detail/limiting_resource_adaptor_impl.cpp b/cpp/src/mr/detail/limiting_resource_adaptor_impl.cpp
index ccc89d6d4..6d7c39ef6 100644
--- a/cpp/src/mr/detail/limiting_resource_adaptor_impl.cpp
+++ b/cpp/src/mr/detail/limiting_resource_adaptor_impl.cpp
@@ -8,6 +8,9 @@
 #include <rmm/detail/format.hpp>
 #include <rmm/mr/detail/limiting_resource_adaptor_impl.hpp>
 
+#include <cuda/stream_ref>
+#include <cuda_runtime_api.h>
+
 namespace RMM_NAMESPACE {
 namespace mr {
 namespace detail {
@@ -69,14 +72,14 @@ void limiting_resource_adaptor_impl::deallocate(cuda::stream_ref stream,
 
 void* limiting_resource_adaptor_impl::allocate_sync(std::size_t bytes, std::size_t alignment)
 {
-  return allocate(cuda_stream_view{}, bytes, alignment);
+  return allocate(cuda::stream_ref{cudaStream_t{nullptr}}, bytes, alignment);
 }
 
 void limiting_resource_adaptor_impl::deallocate_sync(void* ptr,
                                                      std::size_t bytes,
                                                      std::size_t alignment) noexcept
 {
-  deallocate(cuda_stream_view{}, ptr, bytes, alignment);
+  deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
 }
 
 }  // namespace detail
diff --git a/cpp/src/mr/detail/logging_resource_adaptor_impl.cpp b/cpp/src/mr/detail/logging_resource_adaptor_impl.cpp
index 36b8f47e3..749f28e06 100644
--- a/cpp/src/mr/detail/logging_resource_adaptor_impl.cpp
+++ b/cpp/src/mr/detail/logging_resource_adaptor_impl.cpp
@@ -4,10 +4,12 @@
  */
 
 #include <rmm/aligned.hpp>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/detail/format.hpp>
 #include <rmm/mr/detail/logging_resource_adaptor_impl.hpp>
 
+#include <cuda/stream_ref>
+#include <cuda_runtime_api.h>
+
 namespace RMM_NAMESPACE {
 namespace mr {
 namespace detail {
@@ -26,7 +28,7 @@ logging_resource_adaptor_impl::logging_resource_adaptor_impl(
 
 void* logging_resource_adaptor_impl::allocate_sync(std::size_t bytes, std::size_t alignment)
 {
-  auto const stream = cuda_stream_view{};
+  auto const stream = cuda::stream_ref{cudaStream_t{nullptr}};
   try {
     auto const ptr = upstream_mr_.allocate(stream, bytes, alignment);
     logger_->info("allocate,%p,%zu,%s", ptr, bytes, rmm::detail::format_stream(stream));
@@ -41,7 +43,7 @@ void logging_resource_adaptor_impl::deallocate_sync(void* ptr,
                                                     std::size_t bytes,
                                                     std::size_t alignment) noexcept
 {
-  auto const stream = cuda_stream_view{};
+  auto const stream = cuda::stream_ref{cudaStream_t{nullptr}};
   logger_->info("free,%p,%zu,%s", ptr, bytes, rmm::detail::format_stream(stream));
   upstream_mr_.deallocate(stream, ptr, bytes, alignment);
 }
diff --git a/cpp/src/mr/detail/pool_memory_resource_impl.cpp b/cpp/src/mr/detail/pool_memory_resource_impl.cpp
index f00c7fcb8..0ab42c387 100644
--- a/cpp/src/mr/detail/pool_memory_resource_impl.cpp
+++ b/cpp/src/mr/detail/pool_memory_resource_impl.cpp
@@ -4,13 +4,15 @@
  */
 
 #include <rmm/aligned.hpp>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/detail/error.hpp>
 #include <rmm/detail/format.hpp>
 #include <rmm/detail/logging_assert.hpp>
 #include <rmm/logger.hpp>
 #include <rmm/mr/detail/pool_memory_resource_impl.hpp>
 
+#include <cuda/stream_ref>
+#include <cuda_runtime_api.h>
+
 #include <algorithm>
 #include <cstddef>
 #include <limits>
@@ -58,7 +60,7 @@ std::size_t pool_memory_resource_impl::get_maximum_allocation_size() const
 }
 
 pool_memory_resource_impl::block_type pool_memory_resource_impl::try_to_expand(
-  std::size_t try_size, std::size_t min_size, cuda_stream_view stream)
+  std::size_t try_size, std::size_t min_size, cuda::stream_ref stream)
 {
   auto report_error = [&](const char* reason) {
     RMM_LOG_ERROR("[A][Stream %s][Upstream %zuB][FAILURE maximum pool size exceeded: %s]",
@@ -99,13 +101,14 @@ void pool_memory_resource_impl::initialize_pool(std::size_t initial_size,
               "Initial pool size exceeds the maximum pool size!");
 
   if (initial_size > 0) {
-    auto const block = try_to_expand(initial_size, initial_size, cuda_stream_legacy);
-    this->insert_block(block, cuda_stream_legacy);
+    auto const stream = cuda::stream_ref{cudaStreamLegacy};
+    auto const block  = try_to_expand(initial_size, initial_size, stream);
+    this->insert_block(block, stream);
   }
 }
 
 pool_memory_resource_impl::block_type pool_memory_resource_impl::expand_pool(
-  std::size_t size, [[maybe_unused]] free_list& blocks, cuda_stream_view stream)
+  std::size_t size, [[maybe_unused]] free_list& blocks, cuda::stream_ref stream)
 {
   return try_to_expand(size_to_grow(size), size, stream);
 }
@@ -122,7 +125,7 @@ std::size_t pool_memory_resource_impl::size_to_grow(std::size_t size) const
 }
 
 pool_memory_resource_impl::block_type pool_memory_resource_impl::block_from_upstream(
-  std::size_t size, cuda_stream_view stream)
+  std::size_t size, cuda::stream_ref stream)
 {
   RMM_LOG_DEBUG("[A][Stream %s][Upstream %zuB]", rmm::detail::format_stream(stream), size);
 
diff --git a/cpp/src/mr/detail/prefetch_resource_adaptor_impl.cpp b/cpp/src/mr/detail/prefetch_resource_adaptor_impl.cpp
index 31be9763f..1ecd14b04 100644
--- a/cpp/src/mr/detail/prefetch_resource_adaptor_impl.cpp
+++ b/cpp/src/mr/detail/prefetch_resource_adaptor_impl.cpp
@@ -28,7 +28,7 @@ void* prefetch_resource_adaptor_impl::allocate(cuda::stream_ref stream,
                                                std::size_t alignment)
 {
   void* ptr = upstream_mr_.allocate(stream, bytes, alignment);
-  rmm::prefetch(ptr, bytes, rmm::get_current_cuda_device(), cuda_stream_view{stream.get()});
+  rmm::prefetch(ptr, bytes, rmm::get_current_cuda_device(), cuda_stream_view{stream});
   return ptr;
 }
 
@@ -42,14 +42,14 @@ void prefetch_resource_adaptor_impl::deallocate(cuda::stream_ref stream,
 
 void* prefetch_resource_adaptor_impl::allocate_sync(std::size_t bytes, std::size_t alignment)
 {
-  return allocate(cuda_stream_view{}, bytes, alignment);
+  return allocate(cuda::stream_ref{cudaStream_t{nullptr}}, bytes, alignment);
 }
 
 void prefetch_resource_adaptor_impl::deallocate_sync(void* ptr,
                                                      std::size_t bytes,
                                                      std::size_t alignment) noexcept
 {
-  deallocate(cuda_stream_view{}, ptr, bytes, alignment);
+  deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
 }
 
 }  // namespace detail
diff --git a/cpp/src/mr/detail/statistics_resource_adaptor_impl.cpp b/cpp/src/mr/detail/statistics_resource_adaptor_impl.cpp
index 094194c6e..222860713 100644
--- a/cpp/src/mr/detail/statistics_resource_adaptor_impl.cpp
+++ b/cpp/src/mr/detail/statistics_resource_adaptor_impl.cpp
@@ -3,9 +3,11 @@
  * SPDX-License-Identifier: Apache-2.0
  */
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/mr/detail/statistics_resource_adaptor_impl.hpp>
 
+#include <cuda/stream_ref>
+#include <cuda_runtime_api.h>
+
 #include <stdexcept>
 
 namespace RMM_NAMESPACE {
@@ -87,14 +89,14 @@ void statistics_resource_adaptor_impl::deallocate(cuda::stream_ref stream,
 
 void* statistics_resource_adaptor_impl::allocate_sync(std::size_t bytes, std::size_t alignment)
 {
-  return allocate(cuda_stream_view{}, bytes, alignment);
+  return allocate(cuda::stream_ref{cudaStream_t{nullptr}}, bytes, alignment);
 }
 
 void statistics_resource_adaptor_impl::deallocate_sync(void* ptr,
                                                        std::size_t bytes,
                                                        std::size_t alignment) noexcept
 {
-  deallocate(cuda_stream_view{}, ptr, bytes, alignment);
+  deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
 }
 
 }  // namespace detail
diff --git a/cpp/src/mr/detail/thread_safe_resource_adaptor_impl.cpp b/cpp/src/mr/detail/thread_safe_resource_adaptor_impl.cpp
index 91c47e2d4..0a6416818 100644
--- a/cpp/src/mr/detail/thread_safe_resource_adaptor_impl.cpp
+++ b/cpp/src/mr/detail/thread_safe_resource_adaptor_impl.cpp
@@ -5,6 +5,9 @@
 
 #include <rmm/mr/detail/thread_safe_resource_adaptor_impl.hpp>
 
+#include <cuda/stream_ref>
+#include <cuda_runtime_api.h>
+
 namespace RMM_NAMESPACE {
 namespace mr {
 namespace detail {
@@ -40,14 +43,14 @@ void thread_safe_resource_adaptor_impl::deallocate(cuda::stream_ref stream,
 
 void* thread_safe_resource_adaptor_impl::allocate_sync(std::size_t bytes, std::size_t alignment)
 {
-  return allocate(cuda_stream_view{}, bytes, alignment);
+  return allocate(cuda::stream_ref{cudaStream_t{nullptr}}, bytes, alignment);
 }
 
 void thread_safe_resource_adaptor_impl::deallocate_sync(void* ptr,
                                                         std::size_t bytes,
                                                         std::size_t alignment) noexcept
 {
-  deallocate(cuda_stream_view{}, ptr, bytes, alignment);
+  deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
 }
 
 }  // namespace detail
diff --git a/cpp/src/mr/detail/tracking_resource_adaptor_impl.cpp b/cpp/src/mr/detail/tracking_resource_adaptor_impl.cpp
index 5d0ba4231..9cfa95d51 100644
--- a/cpp/src/mr/detail/tracking_resource_adaptor_impl.cpp
+++ b/cpp/src/mr/detail/tracking_resource_adaptor_impl.cpp
@@ -3,10 +3,12 @@
  * SPDX-License-Identifier: Apache-2.0
  */
 
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/logger.hpp>
 #include <rmm/mr/detail/tracking_resource_adaptor_impl.hpp>
 
+#include <cuda/stream_ref>
+#include <cuda_runtime_api.h>
+
 #include <sstream>
 #include <stdexcept>
 
@@ -103,14 +105,14 @@ void tracking_resource_adaptor_impl::deallocate(cuda::stream_ref stream,
 
 void* tracking_resource_adaptor_impl::allocate_sync(std::size_t bytes, std::size_t alignment)
 {
-  return allocate(cuda_stream_view{}, bytes, alignment);
+  return allocate(cuda::stream_ref{cudaStream_t{nullptr}}, bytes, alignment);
 }
 
 void tracking_resource_adaptor_impl::deallocate_sync(void* ptr,
                                                      std::size_t bytes,
                                                      std::size_t alignment) noexcept
 {
-  deallocate(cuda_stream_view{}, ptr, bytes, alignment);
+  deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
 }
 
 }  // namespace detail
diff --git a/cpp/src/prefetch.cpp b/cpp/src/prefetch.cpp
index 0e98291e1..86ae9df42 100644
--- a/cpp/src/prefetch.cpp
+++ b/cpp/src/prefetch.cpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -10,10 +10,11 @@
 
 namespace rmm {
 
-void prefetch(void const* ptr,
-              std::size_t size,
-              rmm::cuda_device_id device,
-              rmm::cuda_stream_view stream)
+namespace {
+void prefetch_impl(void const* ptr,
+                   std::size_t size,
+                   rmm::cuda_device_id device,
+                   cuda::stream_ref stream)
 {
   if (!rmm::detail::concurrent_managed_access::is_supported()) { return; }
 
@@ -22,13 +23,22 @@ void prefetch(void const* ptr,
     (device.value() == cudaCpuDeviceId) ? cudaMemLocationTypeHost : cudaMemLocationTypeDevice,
     device.value()};
   constexpr int flags = 0;
-  cudaError_t result  = cudaMemPrefetchAsync(ptr, size, location, flags, stream.value());
+  cudaError_t result  = cudaMemPrefetchAsync(ptr, size, location, flags, stream.get());
 #else
-  cudaError_t result = cudaMemPrefetchAsync(ptr, size, device.value(), stream.value());
+  cudaError_t result = cudaMemPrefetchAsync(ptr, size, device.value(), stream.get());
 #endif
   // cudaErrorInvalidValue is returned when non-managed memory is passed to
   // cudaMemPrefetchAsync. We treat this as a no-op.
   if (result != cudaErrorInvalidValue && result != cudaSuccess) { RMM_CUDA_TRY(result); }
 }
+}  // namespace
+
+void prefetch(void const* ptr,
+              std::size_t size,
+              rmm::cuda_device_id device,
+              rmm::cuda_stream_view stream)
+{
+  prefetch_impl(ptr, size, device, cuda::stream_ref{stream});
+}
 
 }  // namespace rmm
diff --git a/cpp/tests/container_multidevice_tests.cu b/cpp/tests/container_multidevice_tests.cu
index 0fd81adbf..4fe6cf572 100644
--- a/cpp/tests/container_multidevice_tests.cu
+++ b/cpp/tests/container_multidevice_tests.cu
@@ -39,10 +39,10 @@ TYPED_TEST(ContainerMultiDeviceTest, CreateDestroyDifferentActiveDevice)
 
     {
       if constexpr (std::is_same_v<TypeParam, rmm::device_scalar<int>>) {
-        auto buf = TypeParam(rmm::cuda_stream_view{});
+        auto buf = TypeParam(cuda::stream_ref{cudaStream_t{nullptr}});
         RMM_ASSERT_CUDA_SUCCESS(cudaSetDevice(1));  // force dtor with different active device
       } else {
-        auto buf = TypeParam(128, rmm::cuda_stream_view{});
+        auto buf = TypeParam(128, cuda::stream_ref{cudaStream_t{nullptr}});
         RMM_ASSERT_CUDA_SUCCESS(cudaSetDevice(1));  // force dtor with different active device
       }
     }
@@ -66,19 +66,19 @@ TYPED_TEST(ContainerMultiDeviceTest, CreateMoveDestroyDifferentActiveDevice)
     {
       auto buf_1 = []() {
         if constexpr (std::is_same_v<TypeParam, rmm::device_scalar<int>>) {
-          return TypeParam(rmm::cuda_stream_view{});
+          return TypeParam(cuda::stream_ref{cudaStream_t{nullptr}});
         } else {
-          return TypeParam(128, rmm::cuda_stream_view{});
+          return TypeParam(128, cuda::stream_ref{cudaStream_t{nullptr}});
         }
       }();
 
       {
         if constexpr (std::is_same_v<TypeParam, rmm::device_scalar<int>>) {
           // device_vector does not have a constructor that takes a stream
-          auto buf_0 = TypeParam(rmm::cuda_stream_view{});
+          auto buf_0 = TypeParam(cuda::stream_ref{cudaStream_t{nullptr}});
           buf_1      = std::move(buf_0);
         } else {
-          auto buf_0 = TypeParam(128, rmm::cuda_stream_view{});
+          auto buf_0 = TypeParam(128, cuda::stream_ref{cudaStream_t{nullptr}});
           buf_1      = std::move(buf_0);
         }
       }
@@ -103,9 +103,9 @@ TYPED_TEST(ContainerMultiDeviceTest, ResizeDifferentActiveDevice)
     rmm::mr::set_current_device_resource(device_check_resource_adaptor{orig_mr});
 
     if constexpr (not std::is_same_v<TypeParam, rmm::device_scalar<int>>) {
-      auto buf = TypeParam(128, rmm::cuda_stream_view{});
+      auto buf = TypeParam(128, cuda::stream_ref{cudaStream_t{nullptr}});
       RMM_ASSERT_CUDA_SUCCESS(cudaSetDevice(1));  // force resize with different active device
-      buf.resize(1024, rmm::cuda_stream_view{});
+      buf.resize(1024, cuda::stream_ref{cudaStream_t{nullptr}});
     }
 
     RMM_ASSERT_CUDA_SUCCESS(cudaSetDevice(0));
@@ -125,10 +125,10 @@ TYPED_TEST(ContainerMultiDeviceTest, ShrinkDifferentActiveDevice)
     rmm::mr::set_current_device_resource(device_check_resource_adaptor{orig_mr});
 
     if constexpr (not std::is_same_v<TypeParam, rmm::device_scalar<int>>) {
-      auto buf = TypeParam(128, rmm::cuda_stream_view{});
+      auto buf = TypeParam(128, cuda::stream_ref{cudaStream_t{nullptr}});
       RMM_ASSERT_CUDA_SUCCESS(cudaSetDevice(1));  // force resize with different active device
-      buf.resize(64, rmm::cuda_stream_view{});
-      buf.shrink_to_fit(rmm::cuda_stream_view{});
+      buf.resize(64, cuda::stream_ref{cudaStream_t{nullptr}});
+      buf.shrink_to_fit(cuda::stream_ref{cudaStream_t{nullptr}});
     }
 
     RMM_ASSERT_CUDA_SUCCESS(cudaSetDevice(0));
diff --git a/cpp/tests/cuda_stream_tests.cpp b/cpp/tests/cuda_stream_tests.cpp
index 735a561fc..7a1026db0 100644
--- a/cpp/tests/cuda_stream_tests.cpp
+++ b/cpp/tests/cuda_stream_tests.cpp
@@ -82,7 +82,7 @@ TEST_F(CudaStreamTest, TestStreamViewOstream)
 // Without this we don't get test coverage of ~stream_view, presumably because it is elided
 TEST_F(CudaStreamTest, TestStreamViewDestructor)
 {
-  auto view = std::make_shared<rmm::cuda_stream_view>(rmm::cuda_stream_per_thread);
+  auto view = std::make_shared<rmm::cuda_stream_view>(cuda::stream_ref{cudaStreamPerThread});
   view->synchronize();
 }
 
diff --git a/cpp/tests/device_buffer_tests.cu b/cpp/tests/device_buffer_tests.cu
index eced05302..c21a41ae7 100644
--- a/cpp/tests/device_buffer_tests.cu
+++ b/cpp/tests/device_buffer_tests.cu
@@ -57,26 +57,26 @@ TEST(DeviceBufferSimpleTest, ExplicitResourceRef)
 {
   auto mr = rmm::mr::cuda_memory_resource{};
   rmm::device_async_resource_ref ref{mr};
-  auto buf = rmm::device_buffer(10, rmm::cuda_stream_default, ref);
+  auto buf = rmm::device_buffer(10, cuda::stream_ref{cudaStream_t{nullptr}}, ref);
   EXPECT_EQ(buf.size(), 10);
 }
 
 TYPED_TEST(DeviceBufferTest, EmptyBuffer)
 {
-  rmm::device_buffer buff(0, rmm::cuda_stream_default);
+  rmm::device_buffer buff(0, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_TRUE(buff.is_empty());
 }
 
 TYPED_TEST(DeviceBufferTest, DefaultMemoryResource)
 {
-  rmm::device_buffer buff(this->size, rmm::cuda_stream_default);
+  rmm::device_buffer buff(this->size, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_NE(nullptr, buff.data());
   EXPECT_EQ(this->size, buff.size());
   EXPECT_EQ(this->size, buff.ssize());
   EXPECT_EQ(this->size, buff.capacity());
   EXPECT_EQ(rmm::device_async_resource_ref{rmm::mr::get_current_device_resource_ref()},
             buff.memory_resource());
-  EXPECT_EQ(rmm::cuda_stream_default, buff.stream());
+  EXPECT_EQ(rmm::cuda_stream_view{cudaStream_t{nullptr}}, buff.stream());
 }
 
 TYPED_TEST(DeviceBufferTest, DefaultMemoryResourceStream)
@@ -93,12 +93,12 @@ TYPED_TEST(DeviceBufferTest, DefaultMemoryResourceStream)
 
 TYPED_TEST(DeviceBufferTest, ExplicitMemoryResource)
 {
-  rmm::device_buffer buff(this->size, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer buff(this->size, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
   EXPECT_NE(nullptr, buff.data());
   EXPECT_EQ(this->size, buff.size());
   EXPECT_EQ(this->size, buff.capacity());
   EXPECT_EQ(rmm::device_async_resource_ref{this->mr}, buff.memory_resource());
-  EXPECT_EQ(rmm::cuda_stream_default, buff.stream());
+  EXPECT_EQ(rmm::cuda_stream_view{cudaStream_t{nullptr}}, buff.stream());
 }
 
 TYPED_TEST(DeviceBufferTest, ExplicitMemoryResourceStream)
@@ -116,19 +116,19 @@ TYPED_TEST(DeviceBufferTest, CopyFromRawDevicePointer)
 {
   void* device_memory{nullptr};
   EXPECT_EQ(cudaSuccess, cudaMalloc(&device_memory, this->size));
-  thrust::sequence(rmm::exec_policy(rmm::cuda_stream_default),
+  thrust::sequence(rmm::exec_policy(cuda::stream_ref{cudaStream_t{nullptr}}),
                    static_cast<char*>(device_memory),
                    static_cast<char*>(device_memory) + this->size,
                    0);
-  rmm::device_buffer buff(device_memory, this->size, rmm::cuda_stream_default);
+  rmm::device_buffer buff(device_memory, this->size, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_NE(nullptr, buff.data());
   EXPECT_EQ(this->size, buff.size());
   EXPECT_EQ(this->size, buff.capacity());
   EXPECT_EQ(rmm::device_async_resource_ref{rmm::mr::get_current_device_resource_ref()},
             buff.memory_resource());
-  EXPECT_EQ(rmm::cuda_stream_default, buff.stream());
+  EXPECT_EQ(rmm::cuda_stream_view{cudaStream_t{nullptr}}, buff.stream());
 
-  EXPECT_TRUE(thrust::equal(rmm::exec_policy(rmm::cuda_stream_default),
+  EXPECT_TRUE(thrust::equal(rmm::exec_policy(cuda::stream_ref{cudaStream_t{nullptr}}),
                             static_cast<char*>(device_memory),
                             static_cast<char*>(device_memory) + buff.size(),
                             static_cast<char*>(buff.data())));
@@ -141,13 +141,13 @@ TYPED_TEST(DeviceBufferTest, CopyFromRawHostPointer)
   std::vector<uint8_t> host_data(this->size);
   std::iota(host_data.begin(), host_data.end(), static_cast<uint8_t>(0));
   rmm::device_buffer buff(
-    static_cast<void*>(host_data.data()), this->size, rmm::cuda_stream_default);
+    static_cast<void*>(host_data.data()), this->size, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_NE(nullptr, buff.data());
   EXPECT_EQ(this->size, buff.size());
   EXPECT_EQ(this->size, buff.capacity());
   EXPECT_EQ(rmm::device_async_resource_ref{rmm::mr::get_current_device_resource_ref()},
             buff.memory_resource());
-  EXPECT_EQ(rmm::cuda_stream_default, buff.stream());
+  EXPECT_EQ(rmm::cuda_stream_view{cudaStream_t{nullptr}}, buff.stream());
   buff.stream().synchronize();
   std::vector<uint8_t> host_copy(this->size);
   EXPECT_EQ(cudaSuccess,
@@ -158,42 +158,43 @@ TYPED_TEST(DeviceBufferTest, CopyFromRawHostPointer)
 TYPED_TEST(DeviceBufferTest, CopyFromNullptr)
 {
   // can  copy from a nullptr only if size == 0
-  rmm::device_buffer buff(nullptr, 0, rmm::cuda_stream_default);
+  rmm::device_buffer buff(nullptr, 0, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_EQ(nullptr, buff.data());
   EXPECT_EQ(0, buff.size());
   EXPECT_EQ(0, buff.capacity());
   EXPECT_EQ(rmm::device_async_resource_ref{rmm::mr::get_current_device_resource_ref()},
             buff.memory_resource());
-  EXPECT_EQ(rmm::cuda_stream_default, buff.stream());
+  EXPECT_EQ(rmm::cuda_stream_view{cudaStream_t{nullptr}}, buff.stream());
 }
 
 TYPED_TEST(DeviceBufferTest, CopyFromNullptrNonZero)
 {
   // can  copy from a nullptr only if size == 0
-  EXPECT_THROW(rmm::device_buffer buff(nullptr, 1, rmm::cuda_stream_default), rmm::logic_error);
+  EXPECT_THROW(rmm::device_buffer buff(nullptr, 1, cuda::stream_ref{cudaStream_t{nullptr}}),
+               rmm::logic_error);
 }
 
 TYPED_TEST(DeviceBufferTest, CopyConstructor)
 {
-  rmm::device_buffer buff(this->size, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer buff(this->size, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
 
   // Initialize buffer
-  thrust::sequence(rmm::exec_policy(rmm::cuda_stream_default),
+  thrust::sequence(rmm::exec_policy(cuda::stream_ref{cudaStream_t{nullptr}}),
                    static_cast<char*>(buff.data()),
                    // NOLINTNEXTLINE(cppcoreguidelines-pro-bounds-pointer-arithmetic)
                    static_cast<char*>(buff.data()) + buff.size(),
                    0);
 
-  rmm::device_buffer buff_copy(buff, rmm::cuda_stream_default);  // uses default MR
+  rmm::device_buffer buff_copy(buff, cuda::stream_ref{cudaStream_t{nullptr}});  // uses default MR
   EXPECT_NE(nullptr, buff_copy.data());
   EXPECT_NE(buff.data(), buff_copy.data());
   EXPECT_EQ(buff.size(), buff_copy.size());
   EXPECT_EQ(buff.capacity(), buff_copy.capacity());
   EXPECT_EQ(buff_copy.memory_resource(),
             rmm::device_async_resource_ref{rmm::mr::get_current_device_resource_ref()});
-  EXPECT_EQ(buff_copy.stream(), rmm::cuda_stream_default);
+  EXPECT_EQ(buff_copy.stream(), rmm::cuda_stream_view{cudaStream_t{nullptr}});
 
-  EXPECT_TRUE(thrust::equal(rmm::exec_policy(rmm::cuda_stream_default),
+  EXPECT_TRUE(thrust::equal(rmm::exec_policy(cuda::stream_ref{cudaStream_t{nullptr}}),
                             static_cast<char*>(buff.data()),
                             // NOLINTNEXTLINE(cppcoreguidelines-pro-bounds-pointer-arithmetic)
                             static_cast<char*>(buff.data()) + buff.size(),
@@ -205,7 +206,7 @@ TYPED_TEST(DeviceBufferTest, CopyConstructor)
   EXPECT_EQ(buff_copy2.memory_resource(), buff.memory_resource());
   EXPECT_EQ(buff_copy2.stream(), buff.stream());
 
-  EXPECT_TRUE(thrust::equal(rmm::exec_policy(rmm::cuda_stream_default),
+  EXPECT_TRUE(thrust::equal(rmm::exec_policy(cuda::stream_ref{cudaStream_t{nullptr}}),
                             static_cast<signed char*>(buff.data()),
                             // NOLINTNEXTLINE(cppcoreguidelines-pro-bounds-pointer-arithmetic)
                             static_cast<signed char*>(buff.data()) + buff.size(),
@@ -214,18 +215,18 @@ TYPED_TEST(DeviceBufferTest, CopyConstructor)
 
 TYPED_TEST(DeviceBufferTest, CopyCapacityLargerThanSize)
 {
-  rmm::device_buffer buff(this->size, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer buff(this->size, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
 
   // Resizing smaller to make `size()` < `capacity()`
   auto new_size = this->size - 1;
-  buff.resize(new_size, rmm::cuda_stream_default);
+  buff.resize(new_size, cuda::stream_ref{cudaStream_t{nullptr}});
 
-  thrust::sequence(rmm::exec_policy(rmm::cuda_stream_default),
+  thrust::sequence(rmm::exec_policy(cuda::stream_ref{cudaStream_t{nullptr}}),
                    static_cast<signed char*>(buff.data()),
                    // NOLINTNEXTLINE(cppcoreguidelines-pro-bounds-pointer-arithmetic)
                    static_cast<signed char*>(buff.data()) + buff.size(),
                    0);
-  rmm::device_buffer buff_copy(buff, rmm::cuda_stream_default);
+  rmm::device_buffer buff_copy(buff, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_NE(nullptr, buff_copy.data());
   EXPECT_NE(buff.data(), buff_copy.data());
   EXPECT_EQ(buff.size(), buff_copy.size());
@@ -234,9 +235,9 @@ TYPED_TEST(DeviceBufferTest, CopyCapacityLargerThanSize)
   EXPECT_EQ(new_size, buff_copy.capacity());
   EXPECT_EQ(buff_copy.memory_resource(),
             rmm::device_async_resource_ref{rmm::mr::get_current_device_resource_ref()});
-  EXPECT_EQ(buff_copy.stream(), rmm::cuda_stream_default);
+  EXPECT_EQ(buff_copy.stream(), rmm::cuda_stream_view{cudaStream_t{nullptr}});
 
-  EXPECT_TRUE(thrust::equal(rmm::exec_policy(rmm::cuda_stream_default),
+  EXPECT_TRUE(thrust::equal(rmm::exec_policy(cuda::stream_ref{cudaStream_t{nullptr}}),
                             static_cast<signed char*>(buff.data()),
                             // NOLINTNEXTLINE(cppcoreguidelines-pro-bounds-pointer-arithmetic)
                             static_cast<signed char*>(buff.data()) + buff.size(),
@@ -245,9 +246,9 @@ TYPED_TEST(DeviceBufferTest, CopyCapacityLargerThanSize)
 
 TYPED_TEST(DeviceBufferTest, CopyConstructorExplicitMr)
 {
-  rmm::device_buffer buff(this->size, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer buff(this->size, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
 
-  thrust::sequence(rmm::exec_policy(rmm::cuda_stream_default),
+  thrust::sequence(rmm::exec_policy(cuda::stream_ref{cudaStream_t{nullptr}}),
                    static_cast<signed char*>(buff.data()),
                    // NOLINTNEXTLINE(cppcoreguidelines-pro-bounds-pointer-arithmetic)
                    static_cast<signed char*>(buff.data()) + buff.size(),
@@ -269,13 +270,13 @@ TYPED_TEST(DeviceBufferTest, CopyConstructorExplicitMr)
 
 TYPED_TEST(DeviceBufferTest, CopyCapacityLargerThanSizeExplicitMr)
 {
-  rmm::device_buffer buff(this->size, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer buff(this->size, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
 
   // Resizing smaller to make `size()` < `capacity()`
   auto new_size = this->size - 1;
-  buff.resize(new_size, rmm::cuda_stream_default);
+  buff.resize(new_size, cuda::stream_ref{cudaStream_t{nullptr}});
 
-  thrust::sequence(rmm::exec_policy(rmm::cuda_stream_default),
+  thrust::sequence(rmm::exec_policy(cuda::stream_ref{cudaStream_t{nullptr}}),
                    static_cast<signed char*>(buff.data()),
                    // NOLINTNEXTLINE(cppcoreguidelines-pro-bounds-pointer-arithmetic)
                    static_cast<signed char*>(buff.data()) + buff.size(),
@@ -300,7 +301,7 @@ TYPED_TEST(DeviceBufferTest, CopyCapacityLargerThanSizeExplicitMr)
 
 TYPED_TEST(DeviceBufferTest, MoveConstructor)
 {
-  rmm::device_buffer buff(this->size, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer buff(this->size, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
   auto* ptr     = buff.data();
   auto size     = buff.size();
   auto capacity = buff.capacity();
@@ -321,7 +322,8 @@ TYPED_TEST(DeviceBufferTest, MoveConstructor)
             buff.data());         // NOLINT(bugprone-use-after-move, clang-analyzer-cplusplus.Move)
   EXPECT_EQ(0, buff.size());      // NOLINT(bugprone-use-after-move)
   EXPECT_EQ(0, buff.capacity());  // NOLINT(bugprone-use-after-move)
-  EXPECT_EQ(rmm::cuda_stream_default, buff.stream());  // NOLINT(bugprone-use-after-move)
+  EXPECT_EQ(rmm::cuda_stream_view{cudaStream_t{nullptr}},
+            buff.stream());  // NOLINT(bugprone-use-after-move)
 }
 
 TYPED_TEST(DeviceBufferTest, MoveConstructorStream)
@@ -349,12 +351,13 @@ TYPED_TEST(DeviceBufferTest, MoveConstructorStream)
             buff.data());         // NOLINT(bugprone-use-after-move, clang-analyzer-cplusplus.Move)
   EXPECT_EQ(0, buff.size());      // NOLINT(bugprone-use-after-move)
   EXPECT_EQ(0, buff.capacity());  // NOLINT(bugprone-use-after-move)
-  EXPECT_EQ(rmm::cuda_stream_default, buff.stream());  // NOLINT(bugprone-use-after-move)
+  EXPECT_EQ(rmm::cuda_stream_view{cudaStream_t{nullptr}},
+            buff.stream());  // NOLINT(bugprone-use-after-move)
 }
 
 TYPED_TEST(DeviceBufferTest, MoveAssignmentToDefault)
 {
-  rmm::device_buffer src(this->size, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer src(this->size, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
   auto* ptr     = src.data();
   auto size     = src.size();
   auto capacity = src.capacity();
@@ -376,19 +379,19 @@ TYPED_TEST(DeviceBufferTest, MoveAssignmentToDefault)
   EXPECT_EQ(nullptr, src.data());  // NOLINT(bugprone-use-after-move,clang-analyzer-cplusplus.Move)
   EXPECT_EQ(0, src.size());
   EXPECT_EQ(0, src.capacity());
-  EXPECT_EQ(rmm::cuda_stream_default, src.stream());
+  EXPECT_EQ(rmm::cuda_stream_view{cudaStream_t{nullptr}}, src.stream());
 }
 
 TYPED_TEST(DeviceBufferTest, MoveAssignment)
 {
-  rmm::device_buffer src(this->size, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer src(this->size, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
   auto* ptr     = src.data();
   auto size     = src.size();
   auto capacity = src.capacity();
   auto mr       = src.memory_resource();
   auto stream   = src.stream();
 
-  rmm::device_buffer dest(this->size - 1, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer dest(this->size - 1, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
   dest = std::move(src);
 
   // contents of `from` should be in `to`
@@ -403,12 +406,12 @@ TYPED_TEST(DeviceBufferTest, MoveAssignment)
   EXPECT_EQ(nullptr, src.data());  // NOLINT(bugprone-use-after-move,clang-analyzer-cplusplus.Move)
   EXPECT_EQ(0, src.size());
   EXPECT_EQ(0, src.capacity());
-  EXPECT_EQ(rmm::cuda_stream_default, src.stream());
+  EXPECT_EQ(rmm::cuda_stream_view{cudaStream_t{nullptr}}, src.stream());
 }
 
 TYPED_TEST(DeviceBufferTest, SelfMoveAssignment)
 {
-  rmm::device_buffer buff(this->size, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer buff(this->size, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
   auto* ptr     = buff.data();
   auto size     = buff.size();
   auto capacity = buff.capacity();
@@ -429,9 +432,9 @@ TYPED_TEST(DeviceBufferTest, SelfMoveAssignment)
 
 TYPED_TEST(DeviceBufferTest, ResizeSmaller)
 {
-  rmm::device_buffer buff(this->size, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer buff(this->size, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
 
-  thrust::sequence(rmm::exec_policy(rmm::cuda_stream_default),
+  thrust::sequence(rmm::exec_policy(cuda::stream_ref{cudaStream_t{nullptr}}),
                    static_cast<signed char*>(buff.data()),
                    // NOLINTNEXTLINE(cppcoreguidelines-pro-bounds-pointer-arithmetic)
                    static_cast<signed char*>(buff.data()) + buff.size(),
@@ -439,23 +442,23 @@ TYPED_TEST(DeviceBufferTest, ResizeSmaller)
 
   auto* old_data = buff.data();
   rmm::device_buffer old_content(
-    old_data, buff.size(), rmm::cuda_stream_default, this->mr);  // for comparison
+    old_data, buff.size(), cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);  // for comparison
 
   auto new_size = this->size - 1;
-  buff.resize(new_size, rmm::cuda_stream_default);
+  buff.resize(new_size, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_EQ(new_size, buff.size());
   EXPECT_EQ(this->size, buff.capacity());  // Capacity should be unchanged
   // Resizing smaller means the existing allocation should remain unchanged
   EXPECT_EQ(old_data, buff.data());
 
-  buff.shrink_to_fit(rmm::cuda_stream_default);
+  buff.shrink_to_fit(cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_NE(nullptr, buff.data());
   // A reallocation should have occurred
   EXPECT_NE(old_data, buff.data());
   EXPECT_EQ(new_size, buff.size());
   EXPECT_EQ(buff.capacity(), buff.size());
 
-  EXPECT_TRUE(thrust::equal(rmm::exec_policy(rmm::cuda_stream_default),
+  EXPECT_TRUE(thrust::equal(rmm::exec_policy(cuda::stream_ref{cudaStream_t{nullptr}}),
                             static_cast<signed char*>(buff.data()),
                             // NOLINTNEXTLINE(cppcoreguidelines-pro-bounds-pointer-arithmetic)
                             static_cast<signed char*>(buff.data()) + buff.size(),
@@ -464,10 +467,10 @@ TYPED_TEST(DeviceBufferTest, ResizeSmaller)
 
 TYPED_TEST(DeviceBufferTest, ResizeBigger)
 {
-  rmm::device_buffer buff(this->size, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer buff(this->size, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
   auto* old_data = buff.data();
   auto new_size  = this->size + 1;
-  buff.resize(new_size, rmm::cuda_stream_default);
+  buff.resize(new_size, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_EQ(new_size, buff.size());
   EXPECT_EQ(new_size, buff.capacity());
   // Resizing bigger means the data should point to a new allocation
@@ -476,11 +479,11 @@ TYPED_TEST(DeviceBufferTest, ResizeBigger)
 
 TYPED_TEST(DeviceBufferTest, ReserveSmaller)
 {
-  rmm::device_buffer buff(this->size, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer buff(this->size, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
   auto* const old_data    = buff.data();
   auto const old_capacity = buff.capacity();
   auto const new_capacity = buff.capacity() - 1;
-  buff.reserve(new_capacity, rmm::cuda_stream_default);
+  buff.reserve(new_capacity, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_EQ(this->size, buff.size());
   EXPECT_EQ(old_capacity, buff.capacity());
   // Reserving smaller means the allocation is unchanged
@@ -489,10 +492,10 @@ TYPED_TEST(DeviceBufferTest, ReserveSmaller)
 
 TYPED_TEST(DeviceBufferTest, ReserveBigger)
 {
-  rmm::device_buffer buff(this->size, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer buff(this->size, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
   auto* const old_data    = buff.data();
   auto const new_capacity = buff.capacity() + 1;
-  buff.reserve(new_capacity, rmm::cuda_stream_default);
+  buff.reserve(new_capacity, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_EQ(this->size, buff.size());
   EXPECT_EQ(new_capacity, buff.capacity());
   // Reserving bigger means the data should point to a new allocation
@@ -501,26 +504,27 @@ TYPED_TEST(DeviceBufferTest, ReserveBigger)
 
 TYPED_TEST(DeviceBufferTest, SetGetStream)
 {
-  rmm::device_buffer buff(this->size, rmm::cuda_stream_default, this->mr);
+  rmm::device_buffer buff(this->size, cuda::stream_ref{cudaStream_t{nullptr}}, this->mr);
 
-  EXPECT_EQ(buff.stream(), rmm::cuda_stream_default);
+  EXPECT_EQ(buff.stream(), rmm::cuda_stream_view{cudaStream_t{nullptr}});
 
-  rmm::cuda_stream_view const otherstream{cudaStreamPerThread};
+  auto const otherstream = cuda::stream_ref{cudaStreamPerThread};
   buff.set_stream(otherstream);
 
-  EXPECT_EQ(buff.stream(), otherstream);
+  EXPECT_EQ(buff.stream(), rmm::cuda_stream_view{otherstream});
 }
 
 TEST(DeviceBufferAlignmentTest, DefaultAlignment)
 {
-  rmm::device_buffer buff(100, rmm::cuda_stream_default);
+  rmm::device_buffer buff(100, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_EQ(buff.alignment(), rmm::CUDA_ALLOCATION_ALIGNMENT);
   EXPECT_TRUE(rmm::is_pointer_aligned(buff.data(), rmm::CUDA_ALLOCATION_ALIGNMENT));
 }
 
 TEST(DeviceBufferAlignmentTest, ExplicitAlignmentDefault)
 {
-  rmm::device_buffer buff(100, rmm::CUDA_ALLOCATION_ALIGNMENT, rmm::cuda_stream_default);
+  rmm::device_buffer buff(
+    100, rmm::CUDA_ALLOCATION_ALIGNMENT, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_EQ(buff.alignment(), rmm::CUDA_ALLOCATION_ALIGNMENT);
   EXPECT_TRUE(rmm::is_pointer_aligned(buff.data(), rmm::CUDA_ALLOCATION_ALIGNMENT));
   EXPECT_EQ(buff.size(), 100);
@@ -529,7 +533,7 @@ TEST(DeviceBufferAlignmentTest, ExplicitAlignmentDefault)
 TEST(DeviceBufferAlignmentTest, ExplicitAlignmentSmall)
 {
   constexpr std::size_t alignment{64};
-  rmm::device_buffer buff(100, alignment, rmm::cuda_stream_default);
+  rmm::device_buffer buff(100, alignment, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_EQ(buff.alignment(), alignment);
   EXPECT_TRUE(rmm::is_pointer_aligned(buff.data(), alignment));
   EXPECT_EQ(buff.size(), 100);
@@ -540,13 +544,15 @@ TEST(DeviceBufferAlignmentTest, ExplicitAlignmentSmall)
 TEST(DeviceBufferAlignmentTest, DISABLED_ExplicitAlignmentTooLarge)
 {
   auto constexpr alignment = rmm::CUDA_ALLOCATION_ALIGNMENT * 2;
-  EXPECT_THROW(rmm::device_buffer(100, alignment, rmm::cuda_stream_default), rmm::bad_alloc);
+  EXPECT_THROW(rmm::device_buffer(100, alignment, cuda::stream_ref{cudaStream_t{nullptr}}),
+               rmm::bad_alloc);
 }
 
 TEST(DeviceBufferAlignmentTest, InvalidAlignment)
 {
-  EXPECT_THROW(rmm::device_buffer(100, 0, rmm::cuda_stream_default), rmm::invalid_argument);
-  EXPECT_THROW(std::ignore = rmm::device_buffer(100, 3, rmm::cuda_stream_default),
+  EXPECT_THROW(rmm::device_buffer(100, 0, cuda::stream_ref{cudaStream_t{nullptr}}),
+               rmm::invalid_argument);
+  EXPECT_THROW(std::ignore = rmm::device_buffer(100, 3, cuda::stream_ref{cudaStream_t{nullptr}}),
                rmm::invalid_argument);
 }
 
@@ -554,7 +560,8 @@ TEST(DeviceBufferAlignmentTest, CopyFromSourceExplicitAlignment)
 {
   std::vector<uint8_t> host_data(100, 42);
   std::size_t constexpr alignment{128};
-  rmm::device_buffer buff(host_data.data(), host_data.size(), alignment, rmm::cuda_stream_default);
+  rmm::device_buffer buff(
+    host_data.data(), host_data.size(), alignment, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_EQ(buff.alignment(), alignment);
   EXPECT_TRUE(rmm::is_pointer_aligned(buff.data(), alignment));
   EXPECT_EQ(buff.size(), host_data.size());
@@ -566,24 +573,25 @@ TEST(DeviceBufferAlignmentTest, DISABLED_CopyFromSourceAlignmentTooLarge)
 {
   std::vector<uint8_t> host_data(100, 42);
   auto constexpr alignment = rmm::CUDA_ALLOCATION_ALIGNMENT * 2;
-  EXPECT_THROW(std::ignore = rmm::device_buffer(
-                 host_data.data(), host_data.size(), alignment, rmm::cuda_stream_default),
-               rmm::bad_alloc);
+  EXPECT_THROW(
+    std::ignore = rmm::device_buffer(
+      host_data.data(), host_data.size(), alignment, cuda::stream_ref{cudaStream_t{nullptr}}),
+    rmm::bad_alloc);
 }
 
 TEST(DeviceBufferAlignmentTest, CopyFromSourceInvalidAlignment)
 {
   std::vector<uint8_t> host_data(100, 42);
   EXPECT_THROW(std::ignore = rmm::device_buffer(
-                 host_data.data(), host_data.size(), 3, rmm::cuda_stream_default),
+                 host_data.data(), host_data.size(), 3, cuda::stream_ref{cudaStream_t{nullptr}}),
                rmm::invalid_argument);
 }
 
 TEST(DeviceBufferAlignmentTest, CopyConstructorPreservesAlignment)
 {
   std::size_t constexpr alignment{128};
-  rmm::device_buffer buff(100, alignment, rmm::cuda_stream_default);
-  rmm::device_buffer copy(buff, rmm::cuda_stream_default);
+  rmm::device_buffer buff(100, alignment, cuda::stream_ref{cudaStream_t{nullptr}});
+  rmm::device_buffer copy(buff, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_EQ(copy.alignment(), buff.alignment());
   EXPECT_EQ(copy.alignment(), alignment);
 }
@@ -591,7 +599,7 @@ TEST(DeviceBufferAlignmentTest, CopyConstructorPreservesAlignment)
 TEST(DeviceBufferAlignmentTest, MoveConstructorPreservesAlignment)
 {
   std::size_t constexpr alignment{128};
-  rmm::device_buffer buff(100, alignment, rmm::cuda_stream_default);
+  rmm::device_buffer buff(100, alignment, cuda::stream_ref{cudaStream_t{nullptr}});
   rmm::device_buffer moved(std::move(buff));
   EXPECT_EQ(moved.alignment(), alignment);
   // NOLINTNEXTLINE(bugprone-use-after-move,clang-analyzer-cplusplus.Move)
@@ -601,7 +609,7 @@ TEST(DeviceBufferAlignmentTest, MoveConstructorPreservesAlignment)
 TEST(DeviceBufferAlignmentTest, MoveAssignmentPreservesAlignment)
 {
   std::size_t constexpr alignment{128};
-  rmm::device_buffer src(100, alignment, rmm::cuda_stream_default);
+  rmm::device_buffer src(100, alignment, cuda::stream_ref{cudaStream_t{nullptr}});
   rmm::device_buffer dest;
   dest = std::move(src);
   EXPECT_EQ(dest.alignment(), alignment);
@@ -612,7 +620,7 @@ TEST(DeviceBufferAlignmentTest, MoveAssignmentPreservesAlignment)
 TEST(DeviceBufferAlignmentTest, EmptyBufferAlignment)
 {
   std::size_t constexpr alignment{128};
-  rmm::device_buffer buff(std::size_t{0}, alignment, rmm::cuda_stream_default);
+  rmm::device_buffer buff(std::size_t{0}, alignment, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_EQ(buff.alignment(), alignment);
   EXPECT_EQ(buff.size(), 0);
   EXPECT_TRUE(buff.is_empty());
@@ -622,7 +630,7 @@ TEST(DeviceBufferAlignmentTest, EmptyBufferAlignment)
 TEST(DeviceBufferAlignmentTest, EmptyBufferAlignmentTooLarge)
 {
   auto constexpr alignment = rmm::CUDA_ALLOCATION_ALIGNMENT * 2;
-  rmm::device_buffer buff(std::size_t{0}, alignment, rmm::cuda_stream_default);
+  rmm::device_buffer buff(std::size_t{0}, alignment, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_EQ(buff.alignment(), alignment);
   EXPECT_TRUE(buff.is_empty());
   EXPECT_TRUE(rmm::is_pointer_aligned(buff.data(), alignment));
@@ -637,26 +645,26 @@ TEST(DeviceBufferAlignmentTest, DefaultConstructedHasValidAlignment)
 TEST(DeviceBufferAlignmentTest, DefaultConstructedResizeLarger)
 {
   rmm::device_buffer buff;
-  EXPECT_NO_THROW(buff.resize(100, rmm::cuda_stream_default));
+  EXPECT_NO_THROW(buff.resize(100, cuda::stream_ref{cudaStream_t{nullptr}}));
   EXPECT_EQ(buff.size(), 100);
 }
 
 TEST(DeviceBufferAlignmentTest, DefaultConstructedReserveLarger)
 {
   rmm::device_buffer buff;
-  EXPECT_NO_THROW(buff.reserve(100, rmm::cuda_stream_default));
+  EXPECT_NO_THROW(buff.reserve(100, cuda::stream_ref{cudaStream_t{nullptr}}));
   EXPECT_GE(buff.capacity(), 100);
 }
 
 TEST(DeviceBufferAlignmentTest, DefaultConstructedShrinkToFit)
 {
   rmm::device_buffer buff;
-  EXPECT_NO_THROW(buff.shrink_to_fit(rmm::cuda_stream_default));
+  EXPECT_NO_THROW(buff.shrink_to_fit(cuda::stream_ref{cudaStream_t{nullptr}}));
 }
 
 TEST(DeviceBufferAlignmentTest, EmptyBufferResizeLarger)
 {
-  rmm::device_buffer buff(0, rmm::cuda_stream_default);
-  EXPECT_NO_THROW(buff.resize(100, rmm::cuda_stream_default));
+  rmm::device_buffer buff(0, cuda::stream_ref{cudaStream_t{nullptr}});
+  EXPECT_NO_THROW(buff.resize(100, cuda::stream_ref{cudaStream_t{nullptr}}));
   EXPECT_EQ(buff.size(), 100);
 }
diff --git a/cpp/tests/device_check_resource_adaptor.hpp b/cpp/tests/device_check_resource_adaptor.hpp
index f1724ee30..2080c1560 100644
--- a/cpp/tests/device_check_resource_adaptor.hpp
+++ b/cpp/tests/device_check_resource_adaptor.hpp
@@ -6,7 +6,6 @@
 
 #include <rmm/aligned.hpp>
 #include <rmm/cuda_device.hpp>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/detail/error.hpp>
 #include <rmm/mr/per_device_resource.hpp>
 #include <rmm/resource_ref.hpp>
@@ -57,9 +56,9 @@ class device_check_resource_adaptor final {
 
   void* allocate_sync(std::size_t bytes, std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT)
   {
-    rmm::cuda_stream_view stream{};
-    auto* ptr = allocate(stream, bytes, alignment);
-    stream.synchronize();
+    auto stream = cuda::stream_ref{cudaStream_t{nullptr}};
+    auto* ptr   = allocate(stream, bytes, alignment);
+    RMM_CUDA_TRY(cudaStreamSynchronize(stream.get()));
     return ptr;
   }
 
@@ -67,7 +66,7 @@ class device_check_resource_adaptor final {
                        std::size_t bytes,
                        std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept
   {
-    deallocate(rmm::cuda_stream_view{}, ptr, bytes, alignment);
+    deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
   }
 
   bool operator==(device_check_resource_adaptor const& other) const noexcept
diff --git a/cpp/tests/device_scalar_tests.cpp b/cpp/tests/device_scalar_tests.cpp
index 4f354547b..6f0610dd6 100644
--- a/cpp/tests/device_scalar_tests.cpp
+++ b/cpp/tests/device_scalar_tests.cpp
@@ -137,15 +137,15 @@ TYPED_TEST(DeviceScalarTest, SetGetStream)
 
   EXPECT_EQ(scalar.stream(), this->stream);
 
-  rmm::cuda_stream_view const otherstream{cudaStreamPerThread};
+  auto const otherstream = cuda::stream_ref{cudaStreamPerThread};
   scalar.set_stream(otherstream);
 
-  EXPECT_EQ(scalar.stream(), otherstream);
+  EXPECT_EQ(scalar.stream(), rmm::cuda_stream_view{otherstream});
 }
 
 TEST(DeviceScalarAlignmentTest, SmallAlignment)
 {
-  auto s = rmm::device_scalar<int>(rmm::cuda_stream_view{});
+  auto s = rmm::device_scalar<int>(cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_TRUE(rmm::is_pointer_aligned(s.data(), std::alignment_of_v<decltype(s)::value_type>));
 }
 
@@ -157,6 +157,7 @@ TEST(DeviceScalarAlignmentTest, DISABLED_LargeAlignment)
     int value;
   };
 
-  EXPECT_THROW(std::ignore = rmm::device_scalar<OverAligned>(rmm::cuda_stream_view{}),
-               rmm::bad_alloc);
+  EXPECT_THROW(
+    std::ignore = rmm::device_scalar<OverAligned>(cuda::stream_ref{cudaStream_t{nullptr}}),
+    rmm::bad_alloc);
 }
diff --git a/cpp/tests/device_uvector_tests.cpp b/cpp/tests/device_uvector_tests.cpp
index d4614c304..ce92bb87e 100644
--- a/cpp/tests/device_uvector_tests.cpp
+++ b/cpp/tests/device_uvector_tests.cpp
@@ -27,7 +27,10 @@ template class rmm::device_uvector<int32_t>;
 
 template <typename T>
 struct TypedUVectorTest : ::testing::Test {
-  [[nodiscard]] rmm::cuda_stream_view stream() const noexcept { return rmm::cuda_stream_view{}; }
+  [[nodiscard]] rmm::cuda_stream_view stream() const noexcept
+  {
+    return cuda::stream_ref{cudaStream_t{nullptr}};
+  }
 };
 
 using TestTypes = ::testing::Types<int8_t, int32_t, uint64_t, float, double>;
@@ -244,20 +247,20 @@ TYPED_TEST(TypedUVectorTest, SetElementZeroAsync)
 
 TEST(NegativeZeroTest, PreservesFloatNegativeZero)
 {
-  rmm::device_uvector<float> vec(1, rmm::cuda_stream_view{});
+  rmm::device_uvector<float> vec(1, cuda::stream_ref{cudaStream_t{nullptr}});
   float const neg_zero = -0.0f;
-  vec.set_element_async(0, neg_zero, rmm::cuda_stream_view{});
-  float const result = vec.element(0, rmm::cuda_stream_view{});
+  vec.set_element_async(0, neg_zero, cuda::stream_ref{cudaStream_t{nullptr}});
+  float const result = vec.element(0, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_TRUE(std::signbit(result)) << "sign bit of -0.0f was lost";
   EXPECT_EQ(result, 0.0f);
 }
 
 TEST(NegativeZeroTest, PreservesDoubleNegativeZero)
 {
-  rmm::device_uvector<double> vec(1, rmm::cuda_stream_view{});
+  rmm::device_uvector<double> vec(1, cuda::stream_ref{cudaStream_t{nullptr}});
   double const neg_zero = -0.0;
-  vec.set_element_async(0, neg_zero, rmm::cuda_stream_view{});
-  double const result = vec.element(0, rmm::cuda_stream_view{});
+  vec.set_element_async(0, neg_zero, cuda::stream_ref{cudaStream_t{nullptr}});
+  double const result = vec.element(0, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_TRUE(std::signbit(result)) << "sign bit of -0.0 was lost";
   EXPECT_EQ(result, 0.0);
 }
@@ -283,10 +286,10 @@ TYPED_TEST(TypedUVectorTest, SetGetStream)
 
   EXPECT_EQ(vec.stream(), this->stream());
 
-  rmm::cuda_stream_view const otherstream{cudaStreamPerThread};
+  auto const otherstream = cuda::stream_ref{cudaStreamPerThread};
   vec.set_stream(otherstream);
 
-  EXPECT_EQ(vec.stream(), otherstream);
+  EXPECT_EQ(vec.stream(), rmm::cuda_stream_view{otherstream});
 }
 
 TYPED_TEST(TypedUVectorTest, Iterators)
@@ -383,7 +386,7 @@ TYPED_TEST(TypedUVectorTest, SpanConversionImplicit)
 
 TEST(DeviceUVectorAlignmentTest, SmallAlignment)
 {
-  auto v = rmm::device_uvector<int>(10, rmm::cuda_stream_view{});
+  auto v = rmm::device_uvector<int>(10, cuda::stream_ref{cudaStream_t{nullptr}});
   EXPECT_TRUE(rmm::is_pointer_aligned(v.data(), std::alignment_of_v<decltype(v)::value_type>));
 }
 
@@ -395,6 +398,7 @@ TEST(DeviceUVectorAlignmentTest, DISABLED_LargeAlignment)
     int value;
   };
 
-  EXPECT_THROW(std::ignore = rmm::device_uvector<OverAligned>(10, rmm::cuda_stream_view{}),
-               rmm::bad_alloc);
+  EXPECT_THROW(
+    std::ignore = rmm::device_uvector<OverAligned>(10, cuda::stream_ref{cudaStream_t{nullptr}}),
+    rmm::bad_alloc);
 }
diff --git a/cpp/tests/mock_resource.hpp b/cpp/tests/mock_resource.hpp
index 023da63fe..d8aab9944 100644
--- a/cpp/tests/mock_resource.hpp
+++ b/cpp/tests/mock_resource.hpp
@@ -5,10 +5,10 @@
 #pragma once
 
 #include <rmm/aligned.hpp>
-#include <rmm/cuda_stream_view.hpp>
 
 #include <cuda/memory_resource>
 #include <cuda/stream_ref>
+#include <cuda_runtime_api.h>
 
 #include <gmock/gmock.h>
 
@@ -23,14 +23,14 @@ class mock_resource {
 
   void* allocate_sync(std::size_t bytes, std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT)
   {
-    return allocate(rmm::cuda_stream_view{}, bytes, alignment);
+    return allocate(cuda::stream_ref{cudaStream_t{nullptr}}, bytes, alignment);
   }
 
   void deallocate_sync(void* ptr,
                        std::size_t bytes,
                        std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept
   {
-    deallocate(rmm::cuda_stream_view{}, ptr, bytes, alignment);
+    deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
   }
 
   bool operator==(mock_resource const&) const noexcept { return true; }
diff --git a/cpp/tests/mr/aligned_mr_tests.cpp b/cpp/tests/mr/aligned_mr_tests.cpp
index 9c1da0173..c6bbf9b77 100644
--- a/cpp/tests/mr/aligned_mr_tests.cpp
+++ b/cpp/tests/mr/aligned_mr_tests.cpp
@@ -101,7 +101,7 @@ TEST(AlignedTest, DefaultAllocationAlignmentPassthrough)
   mock_resource_wrapper wrapper{&mock};
   aligned_adaptor mr{device_async_resource_ref{wrapper}};
 
-  cuda_stream_view stream;
+  cuda::stream_ref stream{cudaStream_t{nullptr}};
   void* const pointer = int_to_address(123);
 
   {
@@ -125,7 +125,7 @@ TEST(AlignedTest, BelowAlignmentThresholdPassthrough)
   auto const threshold{65536};
   aligned_adaptor mr{device_async_resource_ref{wrapper}, alignment, threshold};
 
-  cuda_stream_view stream;
+  cuda::stream_ref stream{cudaStream_t{nullptr}};
   void* const pointer = int_to_address(123);
   {
     auto const size{3};
@@ -157,7 +157,7 @@ TEST(AlignedTest, UpstreamAddressAlreadyAligned)
   auto const threshold{65536};
   aligned_adaptor mr{device_async_resource_ref{wrapper}, alignment, threshold};
 
-  cuda_stream_view stream;
+  cuda::stream_ref stream{cudaStream_t{nullptr}};
   void* const pointer = int_to_address(4096);
 
   {
@@ -181,7 +181,7 @@ TEST(AlignedTest, AlignUpstreamAddress)
   auto const threshold{65536};
   aligned_adaptor mr{device_async_resource_ref{wrapper}, alignment, threshold};
 
-  cuda_stream_view stream;
+  cuda::stream_ref stream{cudaStream_t{nullptr}};
   {
     void* const pointer = int_to_address(256);
     auto const size{69376};
@@ -205,7 +205,7 @@ TEST(AlignedTest, AlignMultiple)
   auto const threshold{65536};
   aligned_adaptor mr{device_async_resource_ref{wrapper}, alignment, threshold};
 
-  cuda_stream_view stream;
+  cuda::stream_ref stream{cudaStream_t{nullptr}};
 
   {
     void* const pointer1 = int_to_address(256);
diff --git a/cpp/tests/mr/arena_mr_tests.cpp b/cpp/tests/mr/arena_mr_tests.cpp
index 548bbe76c..88ce1f3be 100644
--- a/cpp/tests/mr/arena_mr_tests.cpp
+++ b/cpp/tests/mr/arena_mr_tests.cpp
@@ -563,8 +563,8 @@ TEST_F(ArenaTest, Defragment)  // NOLINT
     for (std::size_t i = 0; i < num_threads; ++i) {
       threads.emplace_back(std::thread([&] {
         cuda_stream stream{};
-        void* ptr = mr.allocate(cuda_stream_view{stream}, 32_KiB, rmm::CUDA_ALLOCATION_ALIGNMENT);
-        mr.deallocate(cuda_stream_view{stream}, ptr, 32_KiB, rmm::CUDA_ALLOCATION_ALIGNMENT);
+        void* ptr = mr.allocate(cuda::stream_ref{stream}, 32_KiB, rmm::CUDA_ALLOCATION_ALIGNMENT);
+        mr.deallocate(cuda::stream_ref{stream}, ptr, 32_KiB, rmm::CUDA_ALLOCATION_ALIGNMENT);
       }));
     }
     for (auto& thread : threads) {
@@ -585,28 +585,35 @@ TEST_F(ArenaTest, PerThreadToStreamDealloc)  // NOLINT
   auto const arena_size = superblock::minimum_size * 2;
   arena_mr mr(rmm::mr::get_current_device_resource_ref(), arena_size);
   // Create an allocation from a per thread arena
-  void* thread_ptr = mr.allocate(rmm::cuda_stream_per_thread, 256, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  void* thread_ptr =
+    mr.allocate(cuda::stream_ref{cudaStreamPerThread}, 256, rmm::CUDA_ALLOCATION_ALIGNMENT);
   // Create an allocation in a stream arena to force global arena
   // to be empty
   cuda_stream stream{};
-  void* ptr = mr.allocate(cuda_stream_view{stream}, 32_KiB, rmm::CUDA_ALLOCATION_ALIGNMENT);
-  mr.deallocate(cuda_stream_view{stream}, ptr, 32_KiB, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  void* ptr = mr.allocate(cuda::stream_ref{stream}, 32_KiB, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  mr.deallocate(cuda::stream_ref{stream}, ptr, 32_KiB, rmm::CUDA_ALLOCATION_ALIGNMENT);
   // at this point the global arena doesn't have any superblocks so
   // the next allocation causes defrag. Defrag causes all superblocks
   // from the thread and stream arena allocated above to go back to
   // global arena and it allocates one superblock to the stream arena.
-  auto* ptr1 =
-    mr.allocate(rmm::cuda_stream_view{}, superblock::minimum_size, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  auto* ptr1 = mr.allocate(cuda::stream_ref{cudaStream_t{nullptr}},
+                           superblock::minimum_size,
+                           rmm::CUDA_ALLOCATION_ALIGNMENT);
   // Allocate again to make sure all superblocks from
   // global arena are owned by a stream arena instead of a thread arena
   // or the global arena.
-  auto* ptr2 = mr.allocate(rmm::cuda_stream_view{}, 32_KiB, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  auto* ptr2 =
+    mr.allocate(cuda::stream_ref{cudaStream_t{nullptr}}, 32_KiB, rmm::CUDA_ALLOCATION_ALIGNMENT);
   // The original thread ptr is now owned by a stream arena so make
   // sure deallocation works.
-  mr.deallocate(rmm::cuda_stream_per_thread, thread_ptr, 256, rmm::CUDA_ALLOCATION_ALIGNMENT);
   mr.deallocate(
-    rmm::cuda_stream_view{}, ptr1, superblock::minimum_size, rmm::CUDA_ALLOCATION_ALIGNMENT);
-  mr.deallocate(rmm::cuda_stream_view{}, ptr2, 32_KiB, rmm::CUDA_ALLOCATION_ALIGNMENT);
+    cuda::stream_ref{cudaStreamPerThread}, thread_ptr, 256, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  mr.deallocate(cuda::stream_ref{cudaStream_t{nullptr}},
+                ptr1,
+                superblock::minimum_size,
+                rmm::CUDA_ALLOCATION_ALIGNMENT);
+  mr.deallocate(
+    cuda::stream_ref{cudaStream_t{nullptr}}, ptr2, 32_KiB, rmm::CUDA_ALLOCATION_ALIGNMENT);
 }
 
 TEST_F(ArenaTest, DumpLogOnFailure)  // NOLINT
diff --git a/cpp/tests/mr/cccl_mr_ref_test_allocation.hpp b/cpp/tests/mr/cccl_mr_ref_test_allocation.hpp
index a56b02356..99df7d851 100644
--- a/cpp/tests/mr/cccl_mr_ref_test_allocation.hpp
+++ b/cpp/tests/mr/cccl_mr_ref_test_allocation.hpp
@@ -25,7 +25,7 @@ TYPED_TEST_P(CcclMrRefAllocationTest, AllocateDefault) { test_various_allocation
 
 TYPED_TEST_P(CcclMrRefAllocationTest, AllocateDefaultStream)
 {
-  test_various_async_allocations(this->ref, cuda_stream_view{});
+  test_various_async_allocations(this->ref, cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TYPED_TEST_P(CcclMrRefAllocationTest, AllocateOnStream)
@@ -38,7 +38,7 @@ TYPED_TEST_P(CcclMrRefAllocationTest, RandomAllocations) { test_random_allocatio
 TYPED_TEST_P(CcclMrRefAllocationTest, RandomAllocationsDefaultStream)
 {
   test_random_async_allocations(
-    this->ref, default_num_allocations, default_max_size, cuda_stream_view{});
+    this->ref, default_num_allocations, default_max_size, cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TYPED_TEST_P(CcclMrRefAllocationTest, RandomAllocationsStream)
@@ -53,7 +53,8 @@ TYPED_TEST_P(CcclMrRefAllocationTest, MixedRandomAllocationFree)
 
 TYPED_TEST_P(CcclMrRefAllocationTest, MixedRandomAllocationFreeDefaultStream)
 {
-  test_mixed_random_async_allocation_free(this->ref, default_max_size, cuda_stream_view{});
+  test_mixed_random_async_allocation_free(
+    this->ref, default_max_size, cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TYPED_TEST_P(CcclMrRefAllocationTest, MixedRandomAllocationFreeStream)
diff --git a/cpp/tests/mr/cccl_mr_ref_test_basic.hpp b/cpp/tests/mr/cccl_mr_ref_test_basic.hpp
index 2d45e4943..0a5284206 100644
--- a/cpp/tests/mr/cccl_mr_ref_test_basic.hpp
+++ b/cpp/tests/mr/cccl_mr_ref_test_basic.hpp
@@ -27,22 +27,28 @@ TYPED_TEST_P(CcclMrRefTest, SetCurrentDeviceResourceRef)
   auto old = rmm::mr::set_current_device_resource(this->ref);
 
   constexpr std::size_t size{100};
-  void* ptr = old.allocate(rmm::cuda_stream_default, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  void* ptr =
+    old.allocate(cuda::stream_ref{cudaStream_t{nullptr}}, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
   EXPECT_NE(ptr, nullptr);
-  old.deallocate(rmm::cuda_stream_default, ptr, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  old.deallocate(
+    cuda::stream_ref{cudaStream_t{nullptr}}, ptr, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
 
   auto current = rmm::mr::get_current_device_resource_ref();
-  ptr          = current.allocate(rmm::cuda_stream_default, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  ptr =
+    current.allocate(cuda::stream_ref{cudaStream_t{nullptr}}, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
   EXPECT_NE(ptr, nullptr);
-  current.deallocate(rmm::cuda_stream_default, ptr, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  current.deallocate(
+    cuda::stream_ref{cudaStream_t{nullptr}}, ptr, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
 
   test_get_current_device_resource_ref();
 
   rmm::mr::reset_current_device_resource();
   current = rmm::mr::get_current_device_resource_ref();
-  ptr     = current.allocate(rmm::cuda_stream_default, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  ptr =
+    current.allocate(cuda::stream_ref{cudaStream_t{nullptr}}, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
   EXPECT_NE(ptr, nullptr);
-  current.deallocate(rmm::cuda_stream_default, ptr, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  current.deallocate(
+    cuda::stream_ref{cudaStream_t{nullptr}}, ptr, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
 }
 
 TYPED_TEST_P(CcclMrRefTest, SelfEquality) { EXPECT_TRUE(this->ref == this->ref); }
@@ -55,7 +61,7 @@ TYPED_TEST_P(CcclMrRefTest, AllocationsAreDifferent)
 
 TYPED_TEST_P(CcclMrRefTest, AsyncAllocationsAreDifferentDefaultStream)
 {
-  concurrent_async_allocations_are_different(this->ref, cuda_stream_view{});
+  concurrent_async_allocations_are_different(this->ref, cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TYPED_TEST_P(CcclMrRefTest, AsyncAllocationsAreDifferent)
diff --git a/cpp/tests/mr/cccl_mr_ref_test_mt.hpp b/cpp/tests/mr/cccl_mr_ref_test_mt.hpp
index 79ae854d7..771ab11c2 100644
--- a/cpp/tests/mr/cccl_mr_ref_test_mt.hpp
+++ b/cpp/tests/mr/cccl_mr_ref_test_mt.hpp
@@ -82,12 +82,12 @@ TYPED_TEST_P(CcclMrRefTestMT, Allocate)
 
 TYPED_TEST_P(CcclMrRefTestMT, AllocateDefaultStream)
 {
-  spawn(test_various_async_allocations, this->ref, rmm::cuda_stream_view{});
+  spawn(test_various_async_allocations, this->ref, cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TYPED_TEST_P(CcclMrRefTestMT, AllocateOnStream)
 {
-  spawn(test_various_async_allocations, this->ref, this->stream.view());
+  spawn(test_various_async_allocations, this->ref, cuda::stream_ref{this->stream});
 }
 
 TYPED_TEST_P(CcclMrRefTestMT, RandomAllocations)
@@ -101,7 +101,7 @@ TYPED_TEST_P(CcclMrRefTestMT, RandomAllocationsDefaultStream)
         this->ref,
         default_num_allocations,
         default_max_size,
-        rmm::cuda_stream_view{});
+        cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TYPED_TEST_P(CcclMrRefTestMT, RandomAllocationsStream)
@@ -110,7 +110,7 @@ TYPED_TEST_P(CcclMrRefTestMT, RandomAllocationsStream)
         this->ref,
         default_num_allocations,
         default_max_size,
-        this->stream.view());
+        cuda::stream_ref{this->stream});
 }
 
 TYPED_TEST_P(CcclMrRefTestMT, MixedRandomAllocationFree)
@@ -120,36 +120,43 @@ TYPED_TEST_P(CcclMrRefTestMT, MixedRandomAllocationFree)
 
 TYPED_TEST_P(CcclMrRefTestMT, MixedRandomAllocationFreeDefaultStream)
 {
-  spawn(
-    test_mixed_random_async_allocation_free, this->ref, default_max_size, rmm::cuda_stream_view{});
+  spawn(test_mixed_random_async_allocation_free,
+        this->ref,
+        default_max_size,
+        cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TYPED_TEST_P(CcclMrRefTestMT, MixedRandomAllocationFreeStream)
 {
-  spawn(test_mixed_random_async_allocation_free, this->ref, default_max_size, this->stream.view());
+  spawn(test_mixed_random_async_allocation_free,
+        this->ref,
+        default_max_size,
+        cuda::stream_ref{this->stream});
 }
 
 TYPED_TEST_P(CcclMrRefTestMT, AllocFreeDifferentThreadsDefaultStream)
 {
   test_async_allocate_free_different_threads(
-    this->ref, rmm::cuda_stream_default, rmm::cuda_stream_default);
+    this->ref, cuda::stream_ref{cudaStream_t{nullptr}}, cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TYPED_TEST_P(CcclMrRefTestMT, AllocFreeDifferentThreadsPerThreadDefaultStream)
 {
   test_async_allocate_free_different_threads(
-    this->ref, rmm::cuda_stream_per_thread, rmm::cuda_stream_per_thread);
+    this->ref, cuda::stream_ref{cudaStreamPerThread}, cuda::stream_ref{cudaStreamPerThread});
 }
 
 TYPED_TEST_P(CcclMrRefTestMT, AllocFreeDifferentThreadsSameStream)
 {
-  test_async_allocate_free_different_threads(this->ref, this->stream, this->stream);
+  test_async_allocate_free_different_threads(
+    this->ref, cuda::stream_ref{this->stream}, cuda::stream_ref{this->stream});
 }
 
 TYPED_TEST_P(CcclMrRefTestMT, AllocFreeDifferentThreadsDifferentStream)
 {
   rmm::cuda_stream stream_b;
-  test_async_allocate_free_different_threads(this->ref, this->stream, stream_b);
+  test_async_allocate_free_different_threads(
+    this->ref, cuda::stream_ref{this->stream}, cuda::stream_ref{stream_b});
   stream_b.synchronize();
 }
 
diff --git a/cpp/tests/mr/failure_callback_mr_tests.cpp b/cpp/tests/mr/failure_callback_mr_tests.cpp
index bd835680c..26cda1773 100644
--- a/cpp/tests/mr/failure_callback_mr_tests.cpp
+++ b/cpp/tests/mr/failure_callback_mr_tests.cpp
@@ -51,13 +51,13 @@ class always_throw_memory_resource final {
 
   void* allocate_sync(std::size_t bytes, std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT)
   {
-    return allocate(rmm::cuda_stream_view{}, bytes, alignment);
+    return allocate(cuda::stream_ref{cudaStream_t{nullptr}}, bytes, alignment);
   }
   void deallocate_sync(void* ptr,
                        std::size_t bytes,
                        std::size_t alignment = rmm::CUDA_ALLOCATION_ALIGNMENT) noexcept
   {
-    deallocate(rmm::cuda_stream_view{}, ptr, bytes, alignment);
+    deallocate(cuda::stream_ref{cudaStream_t{nullptr}}, ptr, bytes, alignment);
   }
 
   bool operator==(always_throw_memory_resource const&) const noexcept { return true; }
diff --git a/cpp/tests/mr/mr_ref_test.hpp b/cpp/tests/mr/mr_ref_test.hpp
index 4b9ad7360..0101d5e5a 100644
--- a/cpp/tests/mr/mr_ref_test.hpp
+++ b/cpp/tests/mr/mr_ref_test.hpp
@@ -11,7 +11,6 @@
 #include <rmm/aligned.hpp>
 #include <rmm/cuda_device.hpp>
 #include <rmm/cuda_stream.hpp>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/detail/error.hpp>
 #include <rmm/mr/arena_memory_resource.hpp>
 #include <rmm/mr/binning_memory_resource.hpp>
@@ -25,6 +24,8 @@
 #include <rmm/mr/system_memory_resource.hpp>
 #include <rmm/resource_ref.hpp>
 
+#include <cuda/stream_ref>
+
 #include <gtest/gtest.h>
 
 #include <cstddef>
@@ -116,16 +117,16 @@ inline void test_allocate(resource_ref ref, std::size_t bytes)
 
 inline void test_async_allocate(rmm::device_async_resource_ref ref,
                                 std::size_t bytes,
-                                cuda_stream_view stream = {})
+                                cuda::stream_ref stream = cuda::stream_ref{cudaStream_t{nullptr}})
 {
   try {
     void* ptr = ref.allocate(stream, bytes, rmm::CUDA_ALLOCATION_ALIGNMENT);
-    if (not stream.is_default()) { stream.synchronize(); }
+    RMM_CUDA_TRY(cudaStreamSynchronize(stream.get()));
     EXPECT_NE(nullptr, ptr);
     EXPECT_TRUE(is_properly_aligned(ptr));
     EXPECT_TRUE(is_device_accessible_memory(ptr));
     ref.deallocate(stream, ptr, bytes, rmm::CUDA_ALLOCATION_ALIGNMENT);
-    if (not stream.is_default()) { stream.synchronize(); }
+    RMM_CUDA_TRY(cudaStreamSynchronize(stream.get()));
   } catch (rmm::out_of_memory const& e) {
     EXPECT_NE(std::string{e.what()}.find("out_of_memory"), std::string::npos);
   }
@@ -145,7 +146,7 @@ inline void concurrent_allocations_are_different(resource_ref ref)
 }
 
 inline void concurrent_async_allocations_are_different(rmm::device_async_resource_ref ref,
-                                                       cuda_stream_view stream)
+                                                       cuda::stream_ref stream)
 {
   const auto size{8_B};
   void* ptr1 = ref.allocate(stream, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
@@ -187,14 +188,14 @@ inline void test_various_allocations(resource_ref ref)
 }
 
 inline void test_various_async_allocations(rmm::device_async_resource_ref ref,
-                                           cuda_stream_view stream)
+                                           cuda::stream_ref stream)
 {
   // test allocating zero bytes on non-default stream
   {
     void* ptr = ref.allocate(stream, 0, rmm::CUDA_ALLOCATION_ALIGNMENT);
-    stream.synchronize();
+    RMM_CUDA_TRY(cudaStreamSynchronize(stream.get()));
     EXPECT_NO_THROW(ref.deallocate(stream, ptr, 0, rmm::CUDA_ALLOCATION_ALIGNMENT));
-    stream.synchronize();
+    RMM_CUDA_TRY(cudaStreamSynchronize(stream.get()));
   }
 
   test_async_allocate(ref, 4_B, stream);
@@ -244,7 +245,8 @@ inline void test_random_allocations(resource_ref ref,
 inline void test_random_async_allocations(rmm::device_async_resource_ref ref,
                                           std::size_t num_allocations = default_num_allocations,
                                           size_in_bytes max_size      = default_max_size,
-                                          cuda_stream_view stream     = {})
+                                          cuda::stream_ref stream     = cuda::stream_ref{
+                                            cudaStream_t{nullptr}})
 {
   std::vector<allocation> allocations(num_allocations);
 
@@ -258,14 +260,14 @@ inline void test_random_async_allocations(rmm::device_async_resource_ref ref,
     [&generator, &distribution, &ref, stream](allocation& alloc) {
       alloc.size = distribution(generator);
       EXPECT_NO_THROW(alloc.ptr = ref.allocate_sync(alloc.size, rmm::CUDA_ALLOCATION_ALIGNMENT));
-      if (not stream.is_default()) { stream.synchronize(); }
+      RMM_CUDA_TRY(cudaStreamSynchronize(stream.get()));
       EXPECT_NE(nullptr, alloc.ptr);
       EXPECT_TRUE(is_properly_aligned(alloc.ptr));
     });
 
   std::for_each(allocations.begin(), allocations.end(), [stream, &ref](allocation& alloc) {
     EXPECT_NO_THROW(ref.deallocate_sync(alloc.ptr, alloc.size, rmm::CUDA_ALLOCATION_ALIGNMENT));
-    if (not stream.is_default()) { stream.synchronize(); }
+    RMM_CUDA_TRY(cudaStreamSynchronize(stream.get()));
   });
 }
 
@@ -320,7 +322,8 @@ inline void test_mixed_random_allocation_free(resource_ref ref,
 
 inline void test_mixed_random_async_allocation_free(rmm::device_async_resource_ref ref,
                                                     size_in_bytes max_size  = default_max_size,
-                                                    cuda_stream_view stream = {})
+                                                    cuda::stream_ref stream = cuda::stream_ref{
+                                                      cudaStream_t{nullptr}})
 {
   std::default_random_engine generator;
   constexpr std::size_t num_allocations{100};
diff --git a/cpp/tests/mr/mr_ref_test_allocation.hpp b/cpp/tests/mr/mr_ref_test_allocation.hpp
index f3f97c6cd..003c3c1ff 100644
--- a/cpp/tests/mr/mr_ref_test_allocation.hpp
+++ b/cpp/tests/mr/mr_ref_test_allocation.hpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -15,7 +15,7 @@ TEST_P(mr_ref_allocation_test, AllocateDefault) { test_various_allocations(this-
 
 TEST_P(mr_ref_allocation_test, AllocateDefaultStream)
 {
-  test_various_async_allocations(this->ref, cuda_stream_view{});
+  test_various_async_allocations(this->ref, cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TEST_P(mr_ref_allocation_test, AllocateOnStream)
@@ -28,7 +28,7 @@ TEST_P(mr_ref_allocation_test, RandomAllocations) { test_random_allocations(this
 TEST_P(mr_ref_allocation_test, RandomAllocationsDefaultStream)
 {
   test_random_async_allocations(
-    this->ref, default_num_allocations, default_max_size, cuda_stream_view{});
+    this->ref, default_num_allocations, default_max_size, cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TEST_P(mr_ref_allocation_test, RandomAllocationsStream)
@@ -43,7 +43,8 @@ TEST_P(mr_ref_allocation_test, MixedRandomAllocationFree)
 
 TEST_P(mr_ref_allocation_test, MixedRandomAllocationFreeDefaultStream)
 {
-  test_mixed_random_async_allocation_free(this->ref, default_max_size, cuda_stream_view{});
+  test_mixed_random_async_allocation_free(
+    this->ref, default_max_size, cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TEST_P(mr_ref_allocation_test, MixedRandomAllocationFreeStream)
diff --git a/cpp/tests/mr/mr_ref_test_basic.hpp b/cpp/tests/mr/mr_ref_test_basic.hpp
index 3c356f0e4..672e8d9d1 100644
--- a/cpp/tests/mr/mr_ref_test_basic.hpp
+++ b/cpp/tests/mr/mr_ref_test_basic.hpp
@@ -18,15 +18,19 @@ TEST_P(mr_ref_test, SetCurrentDeviceResourceRef)
 
   // Old ref should be functional (verify by successful allocation)
   constexpr std::size_t size{100};
-  void* ptr = old.allocate(rmm::cuda_stream_default, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  void* ptr =
+    old.allocate(cuda::stream_ref{cudaStream_t{nullptr}}, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
   EXPECT_NE(ptr, nullptr);
-  old.deallocate(rmm::cuda_stream_default, ptr, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  old.deallocate(
+    cuda::stream_ref{cudaStream_t{nullptr}}, ptr, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
 
   // Current device resource should be usable for allocation
   auto current = rmm::mr::get_current_device_resource_ref();
-  ptr          = current.allocate(rmm::cuda_stream_default, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  ptr =
+    current.allocate(cuda::stream_ref{cudaStream_t{nullptr}}, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
   EXPECT_NE(ptr, nullptr);
-  current.deallocate(rmm::cuda_stream_default, ptr, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  current.deallocate(
+    cuda::stream_ref{cudaStream_t{nullptr}}, ptr, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
 
   test_get_current_device_resource_ref();
 
@@ -34,9 +38,11 @@ TEST_P(mr_ref_test, SetCurrentDeviceResourceRef)
   rmm::mr::reset_current_device_resource();
   // Verify reset worked by checking allocation succeeds with initial resource
   current = rmm::mr::get_current_device_resource_ref();
-  ptr     = current.allocate(rmm::cuda_stream_default, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  ptr =
+    current.allocate(cuda::stream_ref{cudaStream_t{nullptr}}, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
   EXPECT_NE(ptr, nullptr);
-  current.deallocate(rmm::cuda_stream_default, ptr, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
+  current.deallocate(
+    cuda::stream_ref{cudaStream_t{nullptr}}, ptr, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
 }
 
 TEST_P(mr_ref_test, SelfEquality) { EXPECT_TRUE(this->ref == this->ref); }
@@ -46,7 +52,7 @@ TEST_P(mr_ref_test, AllocationsAreDifferent) { concurrent_allocations_are_differ
 
 TEST_P(mr_ref_test, AsyncAllocationsAreDifferentDefaultStream)
 {
-  concurrent_async_allocations_are_different(this->ref, cuda_stream_view{});
+  concurrent_async_allocations_are_different(this->ref, cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TEST_P(mr_ref_test, AsyncAllocationsAreDifferent)
diff --git a/cpp/tests/mr/mr_ref_test_mt.hpp b/cpp/tests/mr/mr_ref_test_mt.hpp
index 3eec26680..b510a4a02 100644
--- a/cpp/tests/mr/mr_ref_test_mt.hpp
+++ b/cpp/tests/mr/mr_ref_test_mt.hpp
@@ -82,12 +82,12 @@ TEST_P(mr_ref_test_mt, Allocate)
 
 TEST_P(mr_ref_test_mt, AllocateDefaultStream)
 {
-  spawn(test_various_async_allocations, this->ref, rmm::cuda_stream_view{});
+  spawn(test_various_async_allocations, this->ref, cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TEST_P(mr_ref_test_mt, AllocateOnStream)
 {
-  spawn(test_various_async_allocations, this->ref, this->stream.view());
+  spawn(test_various_async_allocations, this->ref, cuda::stream_ref{this->stream});
 }
 
 TEST_P(mr_ref_test_mt, RandomAllocations)
@@ -101,7 +101,7 @@ TEST_P(mr_ref_test_mt, RandomAllocationsDefaultStream)
         this->ref,
         default_num_allocations,
         default_max_size,
-        rmm::cuda_stream_view{});
+        cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TEST_P(mr_ref_test_mt, RandomAllocationsStream)
@@ -110,7 +110,7 @@ TEST_P(mr_ref_test_mt, RandomAllocationsStream)
         this->ref,
         default_num_allocations,
         default_max_size,
-        this->stream.view());
+        cuda::stream_ref{this->stream});
 }
 
 TEST_P(mr_ref_test_mt, MixedRandomAllocationFree)
@@ -120,36 +120,43 @@ TEST_P(mr_ref_test_mt, MixedRandomAllocationFree)
 
 TEST_P(mr_ref_test_mt, MixedRandomAllocationFreeDefaultStream)
 {
-  spawn(
-    test_mixed_random_async_allocation_free, this->ref, default_max_size, rmm::cuda_stream_view{});
+  spawn(test_mixed_random_async_allocation_free,
+        this->ref,
+        default_max_size,
+        cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TEST_P(mr_ref_test_mt, MixedRandomAllocationFreeStream)
 {
-  spawn(test_mixed_random_async_allocation_free, this->ref, default_max_size, this->stream.view());
+  spawn(test_mixed_random_async_allocation_free,
+        this->ref,
+        default_max_size,
+        cuda::stream_ref{this->stream});
 }
 
 TEST_P(mr_ref_test_mt, AllocFreeDifferentThreadsDefaultStream)
 {
   test_async_allocate_free_different_threads(
-    this->ref, rmm::cuda_stream_default, rmm::cuda_stream_default);
+    this->ref, cuda::stream_ref{cudaStream_t{nullptr}}, cuda::stream_ref{cudaStream_t{nullptr}});
 }
 
 TEST_P(mr_ref_test_mt, AllocFreeDifferentThreadsPerThreadDefaultStream)
 {
   test_async_allocate_free_different_threads(
-    this->ref, rmm::cuda_stream_per_thread, rmm::cuda_stream_per_thread);
+    this->ref, cuda::stream_ref{cudaStreamPerThread}, cuda::stream_ref{cudaStreamPerThread});
 }
 
 TEST_P(mr_ref_test_mt, AllocFreeDifferentThreadsSameStream)
 {
-  test_async_allocate_free_different_threads(this->ref, this->stream, this->stream);
+  test_async_allocate_free_different_threads(
+    this->ref, cuda::stream_ref{this->stream}, cuda::stream_ref{this->stream});
 }
 
 TEST_P(mr_ref_test_mt, AllocFreeDifferentThreadsDifferentStream)
 {
   rmm::cuda_stream streamB;
-  test_async_allocate_free_different_threads(this->ref, this->stream, streamB);
+  test_async_allocate_free_different_threads(
+    this->ref, cuda::stream_ref{this->stream}, cuda::stream_ref{streamB});
   streamB.synchronize();
 }
 
diff --git a/cpp/tests/mr/mr_ref_test_mt_helpers.hpp b/cpp/tests/mr/mr_ref_test_mt_helpers.hpp
index b9d19b258..80b999850 100644
--- a/cpp/tests/mr/mr_ref_test_mt_helpers.hpp
+++ b/cpp/tests/mr/mr_ref_test_mt_helpers.hpp
@@ -43,7 +43,7 @@ inline void async_allocate_loop(rmm::device_async_resource_ref ref,
                                 std::mutex& mtx,
                                 std::condition_variable& allocations_ready,
                                 cudaEvent_t& event,
-                                rmm::cuda_stream_view stream)
+                                cuda::stream_ref stream)
 {
   constexpr std::size_t max_size{1_MiB};
 
@@ -55,7 +55,7 @@ inline void async_allocate_loop(rmm::device_async_resource_ref ref,
     void* ptr        = ref.allocate(stream, size, rmm::CUDA_ALLOCATION_ALIGNMENT);
     {
       std::lock_guard<std::mutex> lock(mtx);
-      RMM_CUDA_TRY(cudaEventRecord(event, stream.value()));
+      RMM_CUDA_TRY(cudaEventRecord(event, stream.get()));
       allocations.emplace_back(ptr, size);
     }
     allocations_ready.notify_one();
@@ -71,12 +71,12 @@ inline void async_deallocate_loop(rmm::device_async_resource_ref ref,
                                   std::mutex& mtx,
                                   std::condition_variable& allocations_ready,
                                   cudaEvent_t& event,
-                                  rmm::cuda_stream_view stream)
+                                  cuda::stream_ref stream)
 {
   for (std::size_t i = 0; i < num_allocations; i++) {
     std::unique_lock lock(mtx);
     allocations_ready.wait(lock, [&allocations] { return !allocations.empty(); });
-    RMM_CUDA_TRY(cudaStreamWaitEvent(stream.value(), event));
+    RMM_CUDA_TRY(cudaStreamWaitEvent(stream.get(), event));
     allocation alloc = allocations.front();
     allocations.pop_front();
     ref.deallocate(stream, alloc.ptr, alloc.size, rmm::CUDA_ALLOCATION_ALIGNMENT);
@@ -87,8 +87,8 @@ inline void async_deallocate_loop(rmm::device_async_resource_ref ref,
 }
 
 inline void test_async_allocate_free_different_threads(rmm::device_async_resource_ref ref,
-                                                       rmm::cuda_stream_view streamA,
-                                                       rmm::cuda_stream_view streamB)
+                                                       cuda::stream_ref streamA,
+                                                       cuda::stream_ref streamB)
 {
   constexpr std::size_t num_allocations{100};
 
diff --git a/cpp/tests/mr/pool_mr_tests.cpp b/cpp/tests/mr/pool_mr_tests.cpp
index 42eba3532..3e7f9d7d0 100644
--- a/cpp/tests/mr/pool_mr_tests.cpp
+++ b/cpp/tests/mr/pool_mr_tests.cpp
@@ -89,7 +89,7 @@ TEST(PoolTest, DeletedStream)
   cudaStream_t stream{};  // we don't use rmm::cuda_stream here to make destruction more explicit
   const int size = 10000;
   EXPECT_EQ(cudaSuccess, cudaStreamCreate(&stream));
-  EXPECT_NO_THROW(rmm::device_buffer buff(size, cuda_stream_view{stream}, mr));
+  EXPECT_NO_THROW(rmm::device_buffer buff(size, cuda::stream_ref{stream}, mr));
   EXPECT_EQ(cudaSuccess, cudaStreamDestroy(stream));
   EXPECT_NO_THROW((void)mr.allocate_sync(size));
 }
@@ -151,11 +151,11 @@ TEST(PoolTest, MultidevicePool)
 
     {
       RMM_CUDA_TRY(cudaSetDevice(0));
-      rmm::device_buffer buf_a(16, rmm::cuda_stream_per_thread, mrs[0]);
+      rmm::device_buffer buf_a(16, cuda::stream_ref{cudaStreamPerThread}, mrs[0]);
 
       {
         RMM_CUDA_TRY(cudaSetDevice(1));
-        rmm::device_buffer buf_b(16, rmm::cuda_stream_per_thread, mrs[1]);
+        rmm::device_buffer buf_b(16, cuda::stream_ref{cudaStreamPerThread}, mrs[1]);
       }
 
       RMM_CUDA_TRY(cudaSetDevice(0));
diff --git a/cpp/tests/mr/statistics_mr_tests.cpp b/cpp/tests/mr/statistics_mr_tests.cpp
index 30ae9ac91..20fd11753 100644
--- a/cpp/tests/mr/statistics_mr_tests.cpp
+++ b/cpp/tests/mr/statistics_mr_tests.cpp
@@ -178,7 +178,7 @@ TEST(StatisticsTest, MultiTracking)
   std::vector<std::shared_ptr<rmm::device_buffer>> allocations;
   for (std::size_t i = 0; i < num_allocations; ++i) {
     allocations.emplace_back(
-      std::make_shared<rmm::device_buffer>(ten_MiB, rmm::cuda_stream_default, mr));
+      std::make_shared<rmm::device_buffer>(ten_MiB, cuda::stream_ref{cudaStream_t{nullptr}}, mr));
   }
 
   EXPECT_EQ(mr.get_allocations_counter().value, 10);
@@ -187,8 +187,8 @@ TEST(StatisticsTest, MultiTracking)
 
   rmm::device_async_resource_ref inner_ref{inner_mr};
   for (std::size_t i = 0; i < num_more_allocations; ++i) {
-    allocations.emplace_back(
-      std::make_shared<rmm::device_buffer>(ten_MiB, rmm::cuda_stream_default, inner_ref));
+    allocations.emplace_back(std::make_shared<rmm::device_buffer>(
+      ten_MiB, cuda::stream_ref{cudaStream_t{nullptr}}, inner_ref));
   }
 
   // Check the allocated bytes for both MRs
diff --git a/cpp/tests/mr/thrust_allocator_tests.cu b/cpp/tests/mr/thrust_allocator_tests.cu
index 8cfeffcd7..ddf209621 100644
--- a/cpp/tests/mr/thrust_allocator_tests.cu
+++ b/cpp/tests/mr/thrust_allocator_tests.cu
@@ -6,7 +6,6 @@
 #include "mr_ref_test.hpp"
 
 #include <rmm/cuda_device.hpp>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/detail/error.hpp>
 #include <rmm/device_vector.hpp>
 #include <rmm/mr/per_device_resource.hpp>
@@ -36,8 +35,8 @@ TEST_P(allocator_test, first)
 TEST_P(allocator_test, defaults)
 {
   rmm::mr::set_current_device_resource(this->ref);
-  rmm::mr::thrust_allocator<int> allocator(rmm::cuda_stream_default);
-  EXPECT_EQ(allocator.stream(), rmm::cuda_stream_default);
+  rmm::mr::thrust_allocator<int> allocator(cuda::stream_ref{cudaStream_t{nullptr}});
+  EXPECT_EQ(allocator.stream(), rmm::cuda_stream_view{cudaStream_t{nullptr}});
   EXPECT_EQ(allocator.get_upstream_resource(),
             rmm::device_async_resource_ref{rmm::mr::get_current_device_resource_ref()});
 }
diff --git a/cpp/tests/mr/tracking_mr_tests.cpp b/cpp/tests/mr/tracking_mr_tests.cpp
index ec7b0a6c1..f5188bded 100644
--- a/cpp/tests/mr/tracking_mr_tests.cpp
+++ b/cpp/tests/mr/tracking_mr_tests.cpp
@@ -153,7 +153,7 @@ TEST(TrackingTest, MultiTracking)
   std::vector<std::shared_ptr<rmm::device_buffer>> allocations;
   for (std::size_t i = 0; i < num_allocations; ++i) {
     allocations.emplace_back(
-      std::make_shared<rmm::device_buffer>(ten_MiB, rmm::cuda_stream_default, mr));
+      std::make_shared<rmm::device_buffer>(ten_MiB, cuda::stream_ref{cudaStream_t{nullptr}}, mr));
   }
 
   EXPECT_EQ(mr.get_outstanding_allocations().size(), num_allocations);
@@ -162,8 +162,8 @@ TEST(TrackingTest, MultiTracking)
 
   rmm::device_async_resource_ref inner_ref{inner_mr};
   for (std::size_t i = 0; i < num_more_allocations; ++i) {
-    allocations.emplace_back(
-      std::make_shared<rmm::device_buffer>(ten_MiB, rmm::cuda_stream_default, inner_ref));
+    allocations.emplace_back(std::make_shared<rmm::device_buffer>(
+      ten_MiB, cuda::stream_ref{cudaStream_t{nullptr}}, inner_ref));
   }
 
   // Check the allocated bytes for both MRs