Adopt CUDA stream compatibility accessors (#2557)

## Summary

Use the `get()` and `sync()` compatibility aliases added in [RMM #2537](rapidsai/rmm#2537). These spellings are shared by `rmm::cuda_stream_view` and `cuda::stream_ref`.

This preserves existing stream types and public APIs while extracting mechanical accessor updates from the broader [stream migration](rapidsai/build-planning#318). It is independently buildable without [RMM #2372](rapidsai/rmm#2372) and leaves the migration PR focused on actual type and signature changes.

This updates raw CUDA, library, kernel-launch, and legacy API boundaries throughout cuVS while preserving current stream types. Changes that require RAFT to return `cuda::stream_ref` remain in [cuVS #2521](#2521).

Authors:
  - Bradley Dice (https://github.com/bdice)

Approvers:
  - Divye Gala (https://github.com/divyegala)

URL: #2557

Author: bdice

c/src/core/c_api.cpp

@@ -144,7 +144,7 @@ extern "C" cuvsError_t cuvsStreamGet(cuvsResources_t res, cudaStream_t* stream)
 {
   return cuvs::core::translate_exceptions([=] {
     auto res_ptr = reinterpret_cast<raft::resources*>(res);
-    *stream      = raft::resource::get_cuda_stream(*res_ptr);
+    *stream      = raft::resource::get_cuda_stream(*res_ptr).get();
   });
 }
 

c/src/neighbors/nn_descent.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -107,7 +107,7 @@ void _get_distances(cuvsResources_t res, cuvsNNDescentIndex_t index, DLManagedTe
                     src->data_handle(),
                     dst.extent(0) * dst.extent(1) * sizeof(float),
                     cudaMemcpyDefault,
-                    raft::resource::get_cuda_stream(*res_ptr));
+                    raft::resource::get_cuda_stream(*res_ptr).get());
 
   } else {
     RAFT_FAIL("Unsupported nn-descent index dtype: %d and bits: %d", dtype.code, dtype.bits);

c/tests/neighbors/ann_ivf_sq_c.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -108,7 +108,7 @@ TEST(IvfSqC, BuildSearch)
 
   cuvsResources_t res;
   cuvsResourcesCreate(&res);
-  cuvsStreamSet(res, stream);
+  cuvsStreamSet(res, stream.get());
 
   run_ivf_sq(res,
              n_rows,

c/tests/neighbors/ann_mg_c.cu

@@ -1,10 +1,12 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
 #include <cuda.h>
 #include <gtest/gtest.h>
+#include <cuda/stream>
+
 #include <raft/core/device_mdarray.hpp>
 #include <raft/core/handle.hpp>
 #include <raft/random/rng.cuh>
@@ -158,12 +160,12 @@ class MgCTest : public ::testing::TestWithParam<mg_test_params> {
 
  protected:
   mg_test_params params;
-  rmm::device_uvector<float> index_data{0, rmm::cuda_stream_default};
-  rmm::device_uvector<float> query_data{0, rmm::cuda_stream_default};
-  rmm::device_uvector<int64_t> neighbors_data{0, rmm::cuda_stream_default};
-  rmm::device_uvector<float> distances_data{0, rmm::cuda_stream_default};
-  rmm::device_uvector<int64_t> ref_neighbors_data{0, rmm::cuda_stream_default};
-  rmm::device_uvector<float> ref_distances_data{0, rmm::cuda_stream_default};
+  rmm::device_uvector<float> index_data{0, cuda::stream_ref{cudaStream_t{cudaStreamDefault}}};
+  rmm::device_uvector<float> query_data{0, cuda::stream_ref{cudaStream_t{cudaStreamDefault}}};
+  rmm::device_uvector<int64_t> neighbors_data{0, cuda::stream_ref{cudaStream_t{cudaStreamDefault}}};
+  rmm::device_uvector<float> distances_data{0, cuda::stream_ref{cudaStream_t{cudaStreamDefault}}};
+  rmm::device_uvector<int64_t> ref_neighbors_data{0, cuda::stream_ref{cudaStream_t{cudaStreamDefault}}};
+  rmm::device_uvector<float> ref_distances_data{0, cuda::stream_ref{cudaStream_t{cudaStreamDefault}}};
 
   // Host memory for multi-GPU tests
   std::vector<float> index_data_host;

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

@@ -1,9 +1,10 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
 
+#include <cuda/stream>
 #include <raft/core/error.hpp>
 #include <raft/core/logger_macros.hpp>
 

cpp/include/cuvs/neighbors/common.hpp

@@ -1168,7 +1168,7 @@ auto make_device_dense_row_major_dataset_from_src(raft::resources const& res,
   RAFT_CUDA_TRY(cudaMemsetAsync(out_array.data_handle(),
                                 0,
                                 out_array.size() * sizeof(ValueT),
-                                raft::resource::get_cuda_stream(res)));
+                                raft::resource::get_cuda_stream(res).get()));
   raft::copy_matrix(out_array.data_handle(),
                     target_stride,
                     src.data_handle(),

cpp/internal/cuvs_internal/neighbors/naive_knn.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -11,7 +11,6 @@
 #include <raft/util/cuda_utils.cuh>
 
 #include <raft/core/resource/cuda_stream.hpp>
-#include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/mr/per_device_resource.hpp>
 #include <rmm/resource_ref.hpp>

cpp/src/cluster/detail/agglomerative.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -280,7 +280,7 @@ void extract_flattened_clusters(raft::resources const& handle,
     rmm::device_uvector<value_idx> levels(n_vertices, stream);
 
     value_idx n_blocks = raft::ceildiv(n_vertices, (value_idx)tpb);
-    write_levels_kernel<<<n_blocks, tpb, 0, stream>>>(children, levels.data(), n_vertices);
+    write_levels_kernel<<<n_blocks, tpb, 0, stream.get()>>>(children, levels.data(), n_vertices);
     /**
      * Step 1: Find label roots:
      *
@@ -323,7 +323,7 @@ void extract_flattened_clusters(raft::resources const& handle,
      */
     value_idx cut_level = (n_edges / 2) - (n_clusters - 1);
 
-    inherit_labels<<<n_blocks, tpb, 0, stream>>>(
+    inherit_labels<<<n_blocks, tpb, 0, stream.get()>>>(
       children, levels.data(), n_leaves, tmp_labels.data(), cut_level, n_vertices);
 
     // copy tmp labels to actual labels

cpp/src/cluster/detail/connectivities.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -61,7 +61,7 @@ struct distance_graph_impl<Linkage::KNN_GRAPH, value_idx, value_t> {
     auto thrust_policy = raft::resource::get_thrust_policy(handle);
 
     // Need to symmetrize knn into undirected graph
-    raft::sparse::COO<value_t, value_idx> knn_graph_coo(stream);
+    raft::sparse::COO<value_t, value_idx> knn_graph_coo(stream.get());
 
     auto X_view = raft::make_device_matrix_view<const value_t, value_idx, raft::row_major>(X, m, n);
     cuvs::neighbors::detail::knn_graph<value_idx, value_t, size_t>(
@@ -92,7 +92,7 @@ struct distance_graph_impl<Linkage::KNN_GRAPH, value_idx, value_t> {
       raft::make_const_mdspan(vals_in_view));
 
     raft::sparse::convert::sorted_coo_to_csr(
-      knn_graph_coo.rows(), knn_graph_coo.nnz, indptr.data(), m + 1, stream);
+      knn_graph_coo.rows(), knn_graph_coo.nnz, indptr.data(), m + 1, stream.get());
 
     // TODO: Wouldn't need to copy here if we could compute knn
     // graph directly on the device uvectors
@@ -140,7 +140,7 @@ void pairwise_distances(const raft::resources& handle,
   value_idx nnz = m * m;
 
   value_idx blocks = raft::ceildiv(nnz, (value_idx)256);
-  fill_indices2<value_idx><<<blocks, 256, 0, stream>>>(indices, m, nnz);
+  fill_indices2<value_idx><<<blocks, 256, 0, stream.get()>>>(indices, m, nnz);
 
   raft::linalg::map_offset(handle,
                            raft::make_device_vector_view<value_idx, value_idx>(indptr, m),

cpp/src/cluster/detail/kmeans.cuh

@@ -95,7 +95,7 @@ void kmeansPlusPlus(raft::resources const& handle,
                     rmm::device_uvector<char>& workspace)
 {
   raft::common::nvtx::range<cuvs::common::nvtx::domain::cuvs> fun_scope("kmeansPlusPlus");
-  cudaStream_t stream = raft::resource::get_cuda_stream(handle);
+  cudaStream_t stream = raft::resource::get_cuda_stream(handle).get();
   auto n_samples      = X.extent(0);
   auto n_features     = X.extent(1);
   auto n_clusters     = params.n_clusters;
@@ -309,7 +309,7 @@ void initScalableKMeansPlusPlus(raft::resources const& handle,
 {
   raft::common::nvtx::range<cuvs::common::nvtx::domain::cuvs> fun_scope(
     "initScalableKMeansPlusPlus");
-  cudaStream_t stream = raft::resource::get_cuda_stream(handle);
+  cudaStream_t stream = raft::resource::get_cuda_stream(handle).get();
   auto n_samples      = X.extent(0);
   auto n_features     = X.extent(1);
   auto n_clusters     = params.n_clusters;
@@ -573,7 +573,7 @@ void kmeans_fit(
   auto n_features     = X.extent(1);
   auto n_clusters     = pams.n_clusters;
   auto metric         = pams.metric;
-  cudaStream_t stream = raft::resource::get_cuda_stream(handle);
+  cudaStream_t stream = raft::resource::get_cuda_stream(handle).get();
 
   if (sample_weight.has_value())
     RAFT_EXPECTS(sample_weight.value().extent(0) == n_samples,
@@ -1035,7 +1035,7 @@ void kmeans_predict(raft::resources const& handle,
   raft::common::nvtx::range<cuvs::common::nvtx::domain::cuvs> fun_scope("kmeans_predict");
   auto n_samples      = X.extent(0);
   auto n_features     = X.extent(1);
-  cudaStream_t stream = raft::resource::get_cuda_stream(handle);
+  cudaStream_t stream = raft::resource::get_cuda_stream(handle).get();
   // Check that parameters are valid
   if (sample_weight.has_value())
     RAFT_EXPECTS(sample_weight.value().extent(0) == n_samples,
@@ -1186,7 +1186,7 @@ void kmeans_transform(raft::resources const& handle,
                "kmeans only supports L2Expanded or L2SqrtExpanded distance metrics.");
   raft::common::nvtx::range<cuvs::common::nvtx::domain::cuvs> fun_scope("kmeans_transform");
   raft::default_logger().set_level(pams.verbosity);
-  cudaStream_t stream = raft::resource::get_cuda_stream(handle);
+  cudaStream_t stream = raft::resource::get_cuda_stream(handle).get();
   auto n_samples      = X.extent(0);
   auto n_features     = X.extent(1);
   auto n_clusters     = pams.n_clusters;