Migrate stream APIs from rmm::cuda_stream_view to cuda::stream_ref

c/src/core/c_api.cpp

@@ -15,7 +15,7 @@
 #include <raft/core/resources.hpp>
 #include <raft/util/cudart_utils.hpp>
 #include <rapids_logger/logger.hpp>
-#include <rmm/cuda_stream_view.hpp>
+#include <cuda/stream>
 #include <rmm/mr/cuda_async_memory_resource.hpp>
 #include <rmm/mr/cuda_memory_resource.hpp>
 #include <rmm/mr/managed_memory_resource.hpp>
@@ -136,7 +136,7 @@ extern "C" cuvsError_t cuvsStreamSet(cuvsResources_t res, cudaStream_t stream)
 {
   return cuvs::core::translate_exceptions([=] {
     auto res_ptr = reinterpret_cast<raft::resources*>(res);
-    raft::resource::set_cuda_stream(*res_ptr, static_cast<rmm::cuda_stream_view>(stream));
+    raft::resource::set_cuda_stream(*res_ptr, static_cast<cuda::stream_ref>(stream));
   });
 }
 

cpp/bench/ann/src/common/util.hpp

@@ -171,7 +171,7 @@ inline auto get_stream_from_global_pool() -> cudaStream_t
       detail::global_stream_pool.emplace_back(rmm::cuda_stream::flags::non_blocking);
     }
   }
-  return detail::global_stream_pool[benchmark_thread_id].view();
+  return detail::global_stream_pool[benchmark_thread_id].value();
 #else
   return nullptr;
 #endif

cpp/bench/ann/src/cuvs/cuvs_ann_bench_utils.h

@@ -19,8 +19,8 @@
 #include <raft/core/resource/device_memory_resource.hpp>
 #include <raft/util/cudart_utils.hpp>
 
+#include <cuda/stream>
 #include <rmm/cuda_stream_pool.hpp>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/mr/failure_callback_resource_adaptor.hpp>
 #include <rmm/mr/managed_memory_resource.hpp>
@@ -124,8 +124,8 @@ class configured_raft_resources {
    */
   explicit configured_raft_resources(const std::shared_ptr<shared_raft_resources>& shared_res)
     : shared_res_{shared_res},
-      res_{std::make_unique<raft::device_resources>(
-        rmm::cuda_stream_view(get_stream_from_global_pool()))}
+      res_{
+        std::make_unique<raft::device_resources>(cuda::stream_ref(get_stream_from_global_pool()))}
   {
     raft::resource::set_large_workspace_resource(
       *res_, raft::mr::device_resource{shared_res_->get_large_memory_resource()});

cpp/bench/ann/src/cuvs/cuvs_brute_force_knn.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #include <cuvs/distance/distance.hpp>
@@ -72,7 +72,7 @@ template <typename T, typename DistT = T>
 class BruteForceKNNBenchmark {
  public:
   BruteForceKNNBenchmark(const RandomKNNInputs& params, const std::string& type_str)
-    : stream_(raft::resource::get_cuda_stream(handle_)),
+    : stream_(raft::resource::get_cuda_stream(handle_).get()),
       params_(params),
       type_str_(type_str),
       database(params_.num_db_vecs * params_.dim, stream_),

cpp/bench/ann/src/cuvs/cuvs_cagra_hnswlib_wrapper.h

@@ -45,7 +45,7 @@ class cuvs_cagra_hnswlib : public algo<T>, public algo_gpu {
 
   [[nodiscard]] auto get_sync_stream() const noexcept -> cudaStream_t override
   {
-    return handle_.get_sync_stream();
+    return handle_.get_sync_stream().get();
   }
 
   [[nodiscard]] auto uses_stream() const noexcept -> bool override

cpp/bench/ann/src/cuvs/cuvs_cagra_wrapper.h

@@ -205,7 +205,7 @@ class cuvs_cagra : public algo<T>, public algo_gpu {
 
   [[nodiscard]] auto get_sync_stream() const noexcept -> cudaStream_t override
   {
-    return handle_.get_sync_stream();
+    return handle_.get_sync_stream().get();
   }
 
   [[nodiscard]] auto uses_stream() const noexcept -> bool override

cpp/bench/ann/src/cuvs/cuvs_ivf_flat_wrapper.h

@@ -57,7 +57,7 @@ class cuvs_ivf_flat : public algo<T>, public algo_gpu {
 
   [[nodiscard]] auto get_sync_stream() const noexcept -> cudaStream_t override
   {
-    return handle_.get_sync_stream();
+    return handle_.get_sync_stream().get();
   }
 
   // to enable dataset access from GPU memory

cpp/bench/ann/src/cuvs/cuvs_ivf_pq_wrapper.h

@@ -70,7 +70,7 @@ class cuvs_ivf_pq : public algo<T>, public algo_gpu {
 
   [[nodiscard]] auto get_sync_stream() const noexcept -> cudaStream_t override
   {
-    return handle_.get_sync_stream();
+    return handle_.get_sync_stream().get();
   }
 
   // to enable dataset access from GPU memory

cpp/bench/ann/src/cuvs/cuvs_ivf_rabitq_wrapper.h

@@ -55,7 +55,7 @@ class cuvs_ivf_rabitq : public algo<T>, public algo_gpu {
 
   [[nodiscard]] auto get_sync_stream() const noexcept -> cudaStream_t override
   {
-    return handle_.get_sync_stream();
+    return handle_.get_sync_stream().get();
   }
 
   // to enable dataset access from GPU memory

cpp/bench/ann/src/cuvs/cuvs_ivf_sq_wrapper.h

@@ -52,7 +52,7 @@ class cuvs_ivf_sq : public algo<T>, public algo_gpu {
 
   [[nodiscard]] auto get_sync_stream() const noexcept -> cudaStream_t override
   {
-    return handle_.get_sync_stream();
+    return handle_.get_sync_stream().get();
   }
 
   [[nodiscard]] auto get_preference() const -> algo_property override

cpp/bench/ann/src/cuvs/cuvs_mg_cagra_wrapper.h

@@ -56,7 +56,7 @@ class cuvs_mg_cagra : public algo<T>, public algo_gpu {
   [[nodiscard]] auto get_sync_stream() const noexcept -> cudaStream_t override
   {
     auto stream = raft::resource::get_cuda_stream(clique_);
-    return stream;
+    return stream.get();
   }
 
   // to enable dataset access from GPU memory

cpp/bench/ann/src/cuvs/cuvs_mg_ivf_flat_wrapper.h

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -52,7 +52,7 @@ class cuvs_mg_ivf_flat : public algo<T>, public algo_gpu {
   [[nodiscard]] auto get_sync_stream() const noexcept -> cudaStream_t override
   {
     auto stream = raft::resource::get_cuda_stream(clique_);
-    return stream;
+    return stream.get();
   }
 
   [[nodiscard]] auto uses_stream() const noexcept -> bool override { return false; }

cpp/bench/ann/src/cuvs/cuvs_mg_ivf_pq_wrapper.h

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -52,7 +52,7 @@ class cuvs_mg_ivf_pq : public algo<T>, public algo_gpu {
   [[nodiscard]] auto get_sync_stream() const noexcept -> cudaStream_t override
   {
     auto stream = raft::resource::get_cuda_stream(clique_);
-    return stream;
+    return stream.get();
   }
 
   [[nodiscard]] auto uses_stream() const noexcept -> bool override { return false; }

cpp/bench/ann/src/cuvs/cuvs_vamana_wrapper.h

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -37,7 +37,7 @@ class cuvs_vamana : public algo<T>, public algo_gpu {
 
   [[nodiscard]] auto get_sync_stream() const noexcept -> cudaStream_t override
   {
-    return handle_.get_sync_stream();
+    return handle_.get_sync_stream().get();
   }
 
   // to enable dataset access from GPU memory

cpp/bench/ann/src/cuvs/cuvs_wrapper.h

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023-2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -65,7 +65,7 @@ class cuvs_gpu : public algo<T>, public algo_gpu {
   }
   [[nodiscard]] auto get_sync_stream() const noexcept -> cudaStream_t override
   {
-    return handle_.get_sync_stream();
+    return handle_.get_sync_stream().get();
   }
   void set_search_dataset(const T* dataset, size_t nrow) override;
   void save(const std::string& file) const override;

cpp/cmake/patches/faiss-1.14-cuda-stream-ref.diff

@@ -0,0 +1,68 @@
+diff --git a/faiss/gpu/GpuIndexIVFFlat.cu b/faiss/gpu/GpuIndexIVFFlat.cu
+index aef0b26..d8a6f9c 100644
+--- a/faiss/gpu/GpuIndexIVFFlat.cu
++++ b/faiss/gpu/GpuIndexIVFFlat.cu
+@@ -267,7 +267,7 @@ void GpuIndexIVFFlat::train(idx_t n, const float* x) {
+             // transfer centroids to host
+             auto host_centroids = toHost<float, 2>(
+                     cuvs_ivfflat_index.value().centers().data_handle(),
+-                    raft_handle.get_stream(),
++                    raft_handle.get_stream().get(),
+                     {idx_t(nlist), this->d});
+             quantizer->train(nlist, host_centroids.data());
+             quantizer->add(nlist, host_centroids.data());
+diff --git a/faiss/gpu/GpuIndexIVFPQ.cu b/faiss/gpu/GpuIndexIVFPQ.cu
+index a4d849c..fae7df1 100644
+--- a/faiss/gpu/GpuIndexIVFPQ.cu
++++ b/faiss/gpu/GpuIndexIVFPQ.cu
+@@ -431,7 +431,7 @@ void GpuIndexIVFPQ::train(idx_t n, const float* x) {
+             // transfer centroids to host
+             auto host_centroids = toHost<float, 2>(
+                     cluster_centers.data_handle(),
+-                    raft_handle.get_stream(),
++                    raft_handle.get_stream().get(),
+                     {idx_t(nlist), this->d});
+             quantizer->train(nlist, host_centroids.data());
+             quantizer->add(nlist, host_centroids.data());
+diff --git a/faiss/gpu/impl/CuvsIVFFlat.cu b/faiss/gpu/impl/CuvsIVFFlat.cu
+index 59ccb21..2497891 100644
+--- a/faiss/gpu/impl/CuvsIVFFlat.cu
++++ b/faiss/gpu/impl/CuvsIVFFlat.cu
+@@ -459,7 +459,7 @@ void CuvsIVFFlat::copyInvertedListsFrom(const InvertedLists* ivf) {
+                 cuvs_index->centers().data_handle(),
+                 cuvs_index->dim(),
+                 (uint32_t)nlist,
+-                raft_handle.get_stream());
++                raft_handle.get_stream().get());
+     }
+ }
+
+diff --git a/faiss/gpu/utils/CuvsFilterConvert.cu b/faiss/gpu/utils/CuvsFilterConvert.cu
+index 8a10aee..bbc5f8b 100644
+--- a/faiss/gpu/utils/CuvsFilterConvert.cu
++++ b/faiss/gpu/utils/CuvsFilterConvert.cu
+@@ -91,13 +91,13 @@ void convert_to_bitset_range(
+             (n_elements_to_set + threads_per_block - 1) / threads_per_block;
+
+     if (nbits == original_nbits) {
+-        set_range_kernel<uint32_t><<<blocks, threads_per_block, 0, stream>>>(
++        set_range_kernel<uint32_t><<<blocks, threads_per_block, 0, stream.get()>>>(
+                 (uint32_t*)bitset.data(), imin, imax, n_elements_to_set);
+     } else if (original_nbits == 8) {
+-        set_range_kernel<uint8_t><<<blocks, threads_per_block, 0, stream>>>(
++        set_range_kernel<uint8_t><<<blocks, threads_per_block, 0, stream.get()>>>(
+                 (uint8_t*)bitset.data(), imin, imax, n_elements_to_set);
+     } else if (original_nbits == 64) {
+-        set_range_kernel<uint64_t><<<blocks, threads_per_block, 0, stream>>>(
++        set_range_kernel<uint64_t><<<blocks, threads_per_block, 0, stream.get()>>>(
+                 (uint64_t*)bitset.data(), imin, imax, n_elements_to_set);
+     } else {
+         throw std::invalid_argument("Unsupported original_nbits");
+@@ -175,6 +175,6 @@ void convert_to_bitset_bitmap(
+     const int threads_per_block = 256;
+     const int blocks = (n_elements + threads_per_block - 1) / threads_per_block;
+
+-    set_bitmap_kernel<<<blocks, threads_per_block, 0, stream>>>(
++    set_bitmap_kernel<<<blocks, threads_per_block, 0, stream.get()>>>(
+             bitset.data(), d_bitmap_ptr, n_elements, bitset_original_nbits);
+ }

cpp/cmake/patches/faiss_override.json

@@ -9,6 +9,11 @@
           "file" : "${current_json_dir}/faiss-1.14-cuvs-26.08.diff",
           "issue" : "Multiple fixes for cuVS compatibility. Update Faiss cuVS to be compatible with new Dataset API: update_device_dataset_same_layout now takes dataset_view and make_padded_dataset_view must be called beforehand. Loading an index built from a user-provided KNN graph passes dataset_view into cagra::index, not raw mdspan.",
           "fixed_in" : ""
+        },
+        {
+          "file" : "${current_json_dir}/faiss-1.14-cuda-stream-ref.diff",
+          "issue" : "Unwrap cuda::stream_ref at CUDA kernel launch, Faiss toHost, and legacy RAFT rowNorm boundaries.",
+          "fixed_in" : ""
         }
       ]
     }

cpp/include/cuvs/neighbors/scann.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -17,7 +17,6 @@
 #include <raft/core/resource/stream_view.hpp>
 #include <raft/core/resources.hpp>
 #include <raft/util/integer_utils.hpp>
-#include <rmm/cuda_stream_view.hpp>
 
 #include <cmath>
 #include <cuvs/core/export.hpp>

cpp/include/cuvs/neighbors/vamana.hpp

@@ -16,7 +16,6 @@
 #include <raft/core/resource/stream_view.hpp>
 #include <raft/core/resources.hpp>
 #include <raft/util/integer_utils.hpp>
-#include <rmm/cuda_stream_view.hpp>
 
 #include <cuvs/core/export.hpp>
 #include <optional>

cpp/internal/cuvs_internal/neighbors/refine_helper.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -15,7 +15,7 @@
 #include <raft/core/resources.hpp>
 #include <raft/random/rng.cuh>
 
-#include <rmm/cuda_stream_view.hpp>
+#include <cuda/stream>
 #include <rmm/device_uvector.hpp>
 
 namespace cuvs::neighbors {
@@ -127,7 +127,7 @@ class RefineHelper {
  public:
   RefineInputs<IdxT> p;
   const raft::resources& handle_;
-  rmm::cuda_stream_view stream_;
+  cuda::stream_ref stream_;
 
   raft::device_matrix<DataT, IdxT, row_major> dataset;
   raft::device_matrix<DataT, IdxT, row_major> queries;

cpp/src/cluster/detail/kmeans_balanced.cuh

@@ -898,7 +898,6 @@ void build_clusters(const raft::resources& handle,
                     rmm::device_async_resource_ref device_memory,
                     const MathT* dataset_norm = nullptr)
 {
-  auto stream = raft::resource::get_cuda_stream(handle);
   // "randomly" initialize labels
   auto labels_view = raft::make_device_vector_view<LabelT, IdxT>(cluster_labels, n_rows);
   raft::linalg::map_offset(

cpp/src/neighbors/all_neighbors/all_neighbors_batched.cuh

@@ -131,8 +131,6 @@ void single_gpu_assign_clusters(
   std::optional<raft::device_vector_view<const T, int64_t>> norms_view;
   cuvs::neighbors::brute_force::index<T> brute_force_index(res, centroids, norms_view, metric);
 
-  auto stream = resource::get_cuda_stream(res);
-
   for (size_t i = 0; i < num_batches; i++) {
     size_t row_offset              = n_rows_per_batch * i + base_row_offset;
     size_t n_rows_of_current_batch = std::min(n_rows_per_batch, num_rows - row_offset);