diff --git a/cpp/docs/DEVELOPER_GUIDE.md b/cpp/docs/DEVELOPER_GUIDE.md index 50d4bc63e9..91275558ad 100644 --- a/cpp/docs/DEVELOPER_GUIDE.md +++ b/cpp/docs/DEVELOPER_GUIDE.md @@ -173,7 +173,7 @@ Similar to a `rmm::device_vector`, allocates a contiguous set of elements in dev key differences: - As an optimization, elements are uninitialized and no synchronization occurs at construction. This limits the types `T` to trivially copyable types. -- All operations are stream ordered (i.e., they accept a `cuda_stream_view` specifying the stream +- All operations are stream ordered (i.e., they accept a `cuda::stream_ref` specifying the stream on which the operation is performed). ## Namespaces diff --git a/cpp/examples/developers/graph_operations/graph_operations.cu b/cpp/examples/developers/graph_operations/graph_operations.cu index 9ed8d4bcdc..5180c4c61d 100644 --- a/cpp/examples/developers/graph_operations/graph_operations.cu +++ b/cpp/examples/developers/graph_operations/graph_operations.cu @@ -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 */ @@ -18,6 +18,7 @@ #include #include +#include #include #include @@ -44,7 +45,7 @@ std::unique_ptr initialize_mg_handle() rmm::mr::set_current_device_resource(resource); std::unique_ptr handle = - std::make_unique(rmm::cuda_stream_per_thread, nullptr, resource); + std::make_unique(cuda::stream_ref{cudaStreamPerThread}, nullptr, resource); raft::comms::initialize_mpi_comms(handle.get(), MPI_COMM_WORLD); auto& comm = handle->get_comms(); diff --git a/cpp/examples/developers/vertex_and_edge_partition/vertex_and_edge_partition.cu b/cpp/examples/developers/vertex_and_edge_partition/vertex_and_edge_partition.cu index 5d47daa6e0..d654b41a7a 100644 --- a/cpp/examples/developers/vertex_and_edge_partition/vertex_and_edge_partition.cu +++ b/cpp/examples/developers/vertex_and_edge_partition/vertex_and_edge_partition.cu @@ -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 */ @@ -15,6 +15,7 @@ #include #include +#include #include #include @@ -41,7 +42,7 @@ std::unique_ptr initialize_mg_handle() rmm::mr::set_current_device_resource(resource); std::unique_ptr handle = - std::make_unique(rmm::cuda_stream_per_thread, nullptr, resource); + std::make_unique(cuda::stream_ref{cudaStreamPerThread}, nullptr, resource); raft::comms::initialize_mpi_comms(handle.get(), MPI_COMM_WORLD); auto& comm = handle->get_comms(); diff --git a/cpp/examples/users/multi_gpu_application/mg_graph_algorithms.cpp b/cpp/examples/users/multi_gpu_application/mg_graph_algorithms.cpp index 91bc074960..31195e8c83 100644 --- a/cpp/examples/users/multi_gpu_application/mg_graph_algorithms.cpp +++ b/cpp/examples/users/multi_gpu_application/mg_graph_algorithms.cpp @@ -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 */ @@ -12,6 +12,8 @@ #include #include +#include + #include #include @@ -36,7 +38,7 @@ std::unique_ptr initialize_mg_handle() rmm::mr::set_current_device_resource(resource); std::unique_ptr handle = - std::make_unique(rmm::cuda_stream_per_thread, nullptr, resource); + std::make_unique(cuda::stream_ref{cudaStreamPerThread}, nullptr, resource); raft::comms::initialize_mpi_comms(handle.get(), MPI_COMM_WORLD); auto& comm = handle->get_comms(); diff --git a/cpp/examples/users/single_gpu_application/sg_graph_algorithms.cpp b/cpp/examples/users/single_gpu_application/sg_graph_algorithms.cpp index 77fac22154..35cd6e8fae 100644 --- a/cpp/examples/users/single_gpu_application/sg_graph_algorithms.cpp +++ b/cpp/examples/users/single_gpu_application/sg_graph_algorithms.cpp @@ -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 */ @@ -9,6 +9,8 @@ #include #include +#include + #include #include @@ -19,7 +21,7 @@ std::unique_ptr initialize_sg_handle() rmm::mr::set_current_device_resource(resource); std::unique_ptr handle = - std::make_unique(rmm::cuda_stream_per_thread, nullptr, resource); + std::make_unique(cuda::stream_ref{cudaStreamPerThread}, nullptr, resource); return std::move(handle); } diff --git a/cpp/include/cugraph/dendrogram.hpp b/cpp/include/cugraph/dendrogram.hpp index 0b4a6613f4..8c0cbea354 100644 --- a/cpp/include/cugraph/dendrogram.hpp +++ b/cpp/include/cugraph/dendrogram.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -9,6 +9,8 @@ #include #include +#include + #include #include @@ -19,7 +21,7 @@ class Dendrogram { public: void add_level(vertex_t first_index, vertex_t num_verts, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()) { level_ptr_.push_back( diff --git a/cpp/include/cugraph/detail/decompress_edge_partition.cuh b/cpp/include/cugraph/detail/decompress_edge_partition.cuh index 181b69eba5..6c689d9dd0 100644 --- a/cpp/include/cugraph/detail/decompress_edge_partition.cuh +++ b/cpp/include/cugraph/detail/decompress_edge_partition.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -119,7 +119,7 @@ void decompress_edge_partition_to_fill_edgelist_majors( detail::decompress_to_edgelist_high_degree<<>>( + handle.get_stream().get()>>>( edge_partition, edge_partition.major_range_first(), edge_partition.major_range_first() + (*segment_offsets)[1], @@ -133,7 +133,7 @@ void decompress_edge_partition_to_fill_edgelist_majors( detail::decompress_to_edgelist_mid_degree<<>>( + handle.get_stream().get()>>>( edge_partition, edge_partition.major_range_first() + (*segment_offsets)[1], edge_partition.major_range_first() + (*segment_offsets)[2], diff --git a/cpp/include/cugraph/detail/utility_wrappers.hpp b/cpp/include/cugraph/detail/utility_wrappers.hpp index e8f7be6e3b..75a152eaa7 100644 --- a/cpp/include/cugraph/detail/utility_wrappers.hpp +++ b/cpp/include/cugraph/detail/utility_wrappers.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -12,6 +12,8 @@ #include +#include + namespace CUGRAPH_EXPORT cugraph { namespace detail { @@ -37,7 +39,7 @@ namespace detail { * */ template -void uniform_random_fill(rmm::cuda_stream_view const& stream_view, +void uniform_random_fill(cuda::stream_ref const& stream_view, value_t* d_value, size_t size, value_t min_value, @@ -58,7 +60,7 @@ void uniform_random_fill(rmm::cuda_stream_view const& stream_view, template void transform_increment_ints(raft::device_span values, value_t value, - rmm::cuda_stream_view const& stream_view); + cuda::stream_ref const& stream_view); /** * @ingroup utility_wrappers_cpp @@ -76,7 +78,7 @@ template void transform_not_equal(raft::device_span values, raft::device_span result, value_t compare, - rmm::cuda_stream_view const& stream_view); + cuda::stream_ref const& stream_view); /** * @ingroup utility_wrappers_cpp @@ -94,7 +96,7 @@ void transform_not_equal(raft::device_span values, * @param the maximum value occurring in the edge list */ template -vertex_t compute_maximum_vertex_id(rmm::cuda_stream_view const& stream_view, +vertex_t compute_maximum_vertex_id(cuda::stream_ref const& stream_view, vertex_t const* d_edgelist_srcs, vertex_t const* d_edgelist_dsts, size_t num_edges); @@ -114,7 +116,7 @@ vertex_t compute_maximum_vertex_id(rmm::cuda_stream_view const& stream_view, * @param the maximum value occurring in the edge list */ template -vertex_t compute_maximum_vertex_id(rmm::cuda_stream_view const& stream_view, +vertex_t compute_maximum_vertex_id(cuda::stream_ref const& stream_view, rmm::device_uvector const& d_edgelist_srcs, rmm::device_uvector const& d_edgelist_dsts) { diff --git a/cpp/include/cugraph/dynamic/memory_manager/block_array.hpp b/cpp/include/cugraph/dynamic/memory_manager/block_array.hpp index 8e5fbf1345..3cbfc4883f 100644 --- a/cpp/include/cugraph/dynamic/memory_manager/block_array.hpp +++ b/cpp/include/cugraph/dynamic/memory_manager/block_array.hpp @@ -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 */ #pragma once @@ -8,9 +8,8 @@ #include #include -#include - #include +#include #include #include @@ -31,7 +30,7 @@ class block_array_t { using buffer_type = dataframe_buffer_type_t; - block_array_t(size_t elements_per_block, size_t num_blocks, rmm::cuda_stream_view stream) + block_array_t(size_t elements_per_block, size_t num_blocks, cuda::stream_ref stream) : bit_tree_(elements_per_block, num_blocks), block_storage_(allocate_dataframe_buffer(num_blocks * elements_per_block, stream)) { diff --git a/cpp/include/cugraph/dynamic/memory_manager/block_array_manager.hpp b/cpp/include/cugraph/dynamic/memory_manager/block_array_manager.hpp index a3c7ce2ba6..5d616e27f1 100644 --- a/cpp/include/cugraph/dynamic/memory_manager/block_array_manager.hpp +++ b/cpp/include/cugraph/dynamic/memory_manager/block_array_manager.hpp @@ -1,12 +1,12 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once #include -#include +#include #include #include @@ -55,7 +55,7 @@ class block_array_manager_t { block_array_manager_t(block_array_manager_t&&) = default; block_array_manager_t& operator=(block_array_manager_t&&) = default; - block_access_data_t insert(size_t num_elements_per_block, rmm::cuda_stream_view stream) + block_access_data_t insert(size_t num_elements_per_block, cuda::stream_ref stream) { CUGRAPH_EXPECTS( num_elements_per_block <= max_elements_per_block_array_, diff --git a/cpp/include/cugraph/edge_partition_device_view.cuh b/cpp/include/cugraph/edge_partition_device_view.cuh index 5f6c6fd0ef..7f8759a073 100644 --- a/cpp/include/cugraph/edge_partition_device_view.cuh +++ b/cpp/include/cugraph/edge_partition_device_view.cuh @@ -14,7 +14,6 @@ #include -#include #include #include #include @@ -23,6 +22,7 @@ #include #include #include +#include #include #include #include @@ -169,10 +169,10 @@ __host__ void compute_number_of_edges_with_mask_async_mg( vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (cuda::std::distance(major_first, major_last) == 0) { - RAFT_CUDA_TRY(cudaMemsetAsync(count.data(), 0, sizeof(size_t), stream)); + RAFT_CUDA_TRY(cudaMemsetAsync(count.data(), 0, sizeof(size_t), stream.get())); return; } @@ -190,14 +190,14 @@ __host__ void compute_number_of_edges_with_mask_async_mg( local_degree_first, count.data(), cuda::std::distance(major_first, major_last), - stream); + stream.get()); d_tmp_storage.resize(tmp_storage_bytes, stream); cub::DeviceReduce::Sum(d_tmp_storage.data(), tmp_storage_bytes, local_degree_first, count.data(), cuda::std::distance(major_first, major_last), - stream); + stream.get()); } else { auto local_degree_first = cuda::make_transform_iterator( major_first, @@ -208,14 +208,14 @@ __host__ void compute_number_of_edges_with_mask_async_mg( local_degree_first, count.data(), cuda::std::distance(major_first, major_last), - stream); + stream.get()); d_tmp_storage.resize(tmp_storage_bytes, stream); cub::DeviceReduce::Sum(d_tmp_storage.data(), tmp_storage_bytes, local_degree_first, count.data(), cuda::std::distance(major_first, major_last), - stream); + stream.get()); } } @@ -228,7 +228,7 @@ __host__ void compute_number_of_edges_with_mask_async_mg( vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template __host__ void compute_number_of_edges_with_mask_async_mg( @@ -239,7 +239,7 @@ __host__ void compute_number_of_edges_with_mask_async_mg( vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template __host__ void compute_number_of_edges_with_mask_async_mg( @@ -250,7 +250,7 @@ __host__ void compute_number_of_edges_with_mask_async_mg( vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template __host__ rmm::device_uvector compute_local_degrees_with_mask_mg( @@ -261,7 +261,7 @@ __host__ rmm::device_uvector compute_local_degrees_with_mask_mg( vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { rmm::device_uvector local_degrees(cuda::std::distance(major_first, major_last), stream); if (dcs_nzd_vertices) { @@ -293,7 +293,7 @@ __host__ rmm::device_uvector compute_local_degrees_with_mask_mg( vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template __host__ rmm::device_uvector compute_local_degrees_with_mask_mg( @@ -303,7 +303,7 @@ __host__ rmm::device_uvector compute_local_degrees_with_mask_mg( vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template __host__ void compute_number_of_edges_with_mask_async_sg( @@ -312,10 +312,10 @@ __host__ void compute_number_of_edges_with_mask_async_sg( MajorIterator major_last, raft::device_span count, raft::device_span offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { if (cuda::std::distance(major_first, major_last) == 0) { - RAFT_CUDA_TRY(cudaMemsetAsync(count.data(), 0, sizeof(size_t), stream)); + RAFT_CUDA_TRY(cudaMemsetAsync(count.data(), 0, sizeof(size_t), stream.get())); return; } @@ -334,14 +334,14 @@ __host__ void compute_number_of_edges_with_mask_async_sg( local_degree_first, count.data(), cuda::std::distance(major_first, major_last), - stream); + stream.get()); d_tmp_storage.resize(tmp_storage_bytes, stream); cub::DeviceReduce::Sum(d_tmp_storage.data(), tmp_storage_bytes, local_degree_first, count.data(), cuda::std::distance(major_first, major_last), - stream); + stream.get()); } template @@ -350,7 +350,7 @@ __host__ void compute_number_of_edges_with_mask_async_sg( raft::device_span majors, raft::device_span count, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template __host__ void compute_number_of_edges_with_mask_async_sg( @@ -358,7 +358,7 @@ __host__ void compute_number_of_edges_with_mask_async_sg( std::tuple vertex_partition_range, raft::device_span count, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template __host__ rmm::device_uvector compute_local_degrees_with_mask_sg( @@ -366,7 +366,7 @@ __host__ rmm::device_uvector compute_local_degrees_with_mask_sg( MajorIterator major_first, MajorIterator major_last, raft::device_span offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { rmm::device_uvector local_degrees(cuda::std::distance(major_first, major_last), stream); thrust::transform( @@ -387,14 +387,14 @@ __host__ rmm::device_uvector compute_local_degrees_with_mask_sg( cuda::std::optional edge_mask, raft::device_span majors, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template __host__ rmm::device_uvector compute_local_degrees_with_mask_sg( cuda::std::optional edge_mask, std::tuple vertex_partition_range, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); } // namespace detail @@ -423,7 +423,7 @@ class edge_partition_device_view_t edge_mask, raft::device_span majors, raft::device_span count, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { detail::compute_number_of_edges_with_mask_async_mg( cuda::std::optional{edge_mask.data()}, @@ -440,7 +440,7 @@ class edge_partition_device_view_t edge_mask, std::tuple vertex_partition_range, raft::device_span count, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { detail::compute_number_of_edges_with_mask_async_mg( cuda::std::optional{edge_mask.data()}, @@ -458,7 +458,7 @@ class edge_partition_device_view_t count, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { detail::compute_number_of_edges_with_mask_async_mg( cuda::std::optional{edge_mask.data()}, @@ -478,7 +478,7 @@ class edge_partition_device_view_t edge_mask, majors_from_offsets_t majors, raft::device_span count, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { detail::compute_number_of_edges_with_mask_async_mg( cuda::std::optional{edge_mask.data()}, @@ -493,7 +493,7 @@ class edge_partition_device_view_t majors, raft::device_span count, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { detail::compute_number_of_edges_with_mask_async_mg(cuda::std::nullopt, majors, @@ -508,7 +508,7 @@ class edge_partition_device_view_t __host__ void compute_number_of_edges_async(raft::device_span majors, raft::device_span count, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { detail::compute_number_of_edges_with_mask_async_mg(cuda::std::nullopt, majors, @@ -522,7 +522,7 @@ class edge_partition_device_view_t vertex_partition_range, raft::device_span count, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { detail::compute_number_of_edges_with_mask_async_mg(cuda::std::nullopt, vertex_partition_range, @@ -539,7 +539,7 @@ class edge_partition_device_view_t count, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { detail::compute_number_of_edges_with_mask_async_mg(cuda::std::nullopt, major_first, @@ -555,7 +555,7 @@ class edge_partition_device_view_t __host__ size_t compute_number_of_edges_with_mask(raft::device_span edge_mask, raft::device_span majors, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { rmm::device_scalar count(stream); count.set_value_to_zero_async(stream); @@ -574,7 +574,7 @@ class edge_partition_device_view_t edge_mask, std::tuple vertex_partition_range, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { rmm::device_scalar count(stream); count.set_value_to_zero_async(stream); @@ -593,7 +593,7 @@ class edge_partition_device_view_t edge_mask, majors_from_offsets_t majors, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { rmm::device_scalar count(stream); count.set_value_to_zero_async(stream); @@ -614,7 +614,7 @@ class edge_partition_device_view_t edge_mask, MajorIterator major_first, MajorIterator major_last, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { rmm::device_scalar count(stream); count.set_value_to_zero_async(stream); @@ -633,7 +633,7 @@ class edge_partition_device_view_t __host__ size_t compute_number_of_edges(raft::device_span majors, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { rmm::device_scalar count(stream); count.set_value_to_zero_async(stream); @@ -642,7 +642,7 @@ class edge_partition_device_view_t vertex_partition_range, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { rmm::device_scalar count(stream); count.set_value_to_zero_async(stream); @@ -652,7 +652,7 @@ class edge_partition_device_view_t majors, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { rmm::device_scalar count(stream); count.set_value_to_zero_async(stream); @@ -671,7 +671,7 @@ class edge_partition_device_view_t, int> = 0> __host__ size_t compute_number_of_edges(MajorIterator major_first, MajorIterator major_last, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { rmm::device_scalar count(stream); count.set_value_to_zero_async(stream); @@ -681,7 +681,7 @@ class edge_partition_device_view_t compute_local_degrees_with_mask( - raft::device_span edge_mask, rmm::cuda_stream_view stream) const + raft::device_span edge_mask, cuda::stream_ref stream) const { return detail::compute_local_degrees_with_mask_mg( cuda::std::optional{edge_mask.data()}, @@ -697,7 +697,7 @@ class edge_partition_device_view_t compute_local_degrees_with_mask( raft::device_span edge_mask, raft::device_span majors, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { return detail::compute_local_degrees_with_mask_mg( cuda::std::optional{edge_mask.data()}, @@ -712,7 +712,7 @@ class edge_partition_device_view_t compute_local_degrees_with_mask( raft::device_span edge_mask, std::tuple vertex_partition_range, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { return detail::compute_local_degrees_with_mask_mg( cuda::std::optional{edge_mask.data()}, @@ -730,7 +730,7 @@ class edge_partition_device_view_t edge_mask, MajorIterator major_first, MajorIterator major_last, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { return detail::compute_local_degrees_with_mask_mg( cuda::std::optional{edge_mask.data()}, @@ -745,7 +745,7 @@ class edge_partition_device_view_t __host__ rmm::device_uvector compute_local_degrees(raft::device_span majors, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { return detail::compute_local_degrees_with_mask_mg(cuda::std::nullopt, majors, @@ -757,7 +757,7 @@ class edge_partition_device_view_t compute_local_degrees( - std::tuple vertex_partition_range, rmm::cuda_stream_view stream) const + std::tuple vertex_partition_range, cuda::stream_ref stream) const { return detail::compute_local_degrees_with_mask_mg(cuda::std::nullopt, vertex_partition_range, @@ -772,7 +772,7 @@ class edge_partition_device_view_t, int> = 0> __host__ rmm::device_uvector compute_local_degrees(MajorIterator major_first, MajorIterator major_last, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { return detail::compute_local_degrees_with_mask_mg(cuda::std::nullopt, major_first, @@ -930,7 +930,7 @@ class edge_partition_device_view_t edge_mask, raft::device_span majors, raft::device_span count, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { detail::compute_number_of_edges_with_mask_async_sg( cuda::std::optional{edge_mask.data()}, @@ -944,7 +944,7 @@ class edge_partition_device_view_t edge_mask, std::tuple vertex_partition_range, raft::device_span count, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { detail::compute_number_of_edges_with_mask_async_sg( cuda::std::optional{edge_mask.data()}, @@ -960,7 +960,7 @@ class edge_partition_device_view_t count, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { detail::compute_number_of_edges_with_mask_async_sg( cuda::std::optional{edge_mask.data()}, @@ -974,7 +974,7 @@ class edge_partition_device_view_t __host__ void compute_number_of_edges_async(raft::device_span majors, raft::device_span count, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { detail::compute_number_of_edges_with_mask_async_sg( cuda::std::nullopt, majors, count, this->offsets_, stream); @@ -982,7 +982,7 @@ class edge_partition_device_view_t vertex_partition_range, raft::device_span count, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { detail::compute_number_of_edges_with_mask_async_sg( cuda::std::nullopt, vertex_partition_range, count, this->offsets_, stream); @@ -993,7 +993,7 @@ class edge_partition_device_view_t count, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { detail::compute_number_of_edges_with_mask_async_sg( cuda::std::nullopt, major_first, major_last, count, this->offsets_, stream); @@ -1002,7 +1002,7 @@ class edge_partition_device_view_t __host__ size_t compute_number_of_edges_with_mask(raft::device_span edge_mask, raft::device_span majors, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { rmm::device_scalar count(stream); count.set_value_to_zero_async(stream); @@ -1014,7 +1014,7 @@ class edge_partition_device_view_t edge_mask, std::tuple vertex_partition_range, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { rmm::device_scalar count(stream); count.set_value_to_zero_async(stream); @@ -1028,7 +1028,7 @@ class edge_partition_device_view_t edge_mask, MajorIterator major_first, MajorIterator major_last, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { rmm::device_scalar count(stream); count.set_value_to_zero_async(stream); @@ -1039,7 +1039,7 @@ class edge_partition_device_view_t __host__ size_t compute_number_of_edges(raft::device_span majors, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { rmm::device_scalar count(stream); count.set_value_to_zero_async(stream); @@ -1048,7 +1048,7 @@ class edge_partition_device_view_t vertex_partition_range, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { rmm::device_scalar count(stream); count.set_value_to_zero_async(stream); @@ -1061,7 +1061,7 @@ class edge_partition_device_view_t, int> = 0> __host__ size_t compute_number_of_edges(MajorIterator major_first, MajorIterator major_last, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { rmm::device_scalar count(stream); count.set_value_to_zero_async(stream); @@ -1071,7 +1071,7 @@ class edge_partition_device_view_t compute_local_degrees_with_mask( - raft::device_span edge_mask, rmm::cuda_stream_view stream) const + raft::device_span edge_mask, cuda::stream_ref stream) const { return detail::compute_local_degrees_with_mask_sg( cuda::std::optional{edge_mask.data()}, @@ -1084,7 +1084,7 @@ class edge_partition_device_view_t compute_local_degrees_with_mask( raft::device_span edge_mask, raft::device_span majors, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { return detail::compute_local_degrees_with_mask_sg( cuda::std::optional{edge_mask.data()}, majors, this->offsets_, stream); @@ -1093,7 +1093,7 @@ class edge_partition_device_view_t compute_local_degrees_with_mask( raft::device_span edge_mask, std::tuple vertex_partition_range, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { return detail::compute_local_degrees_with_mask_sg( cuda::std::optional{edge_mask.data()}, @@ -1108,7 +1108,7 @@ class edge_partition_device_view_t edge_mask, MajorIterator major_first, MajorIterator major_last, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { return detail::compute_local_degrees_with_mask_sg( cuda::std::optional{edge_mask.data()}, @@ -1120,14 +1120,14 @@ class edge_partition_device_view_t __host__ rmm::device_uvector compute_local_degrees(raft::device_span majors, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { return detail::compute_local_degrees_with_mask_sg( cuda::std::nullopt, majors, this->offsets_, stream); } __host__ rmm::device_uvector compute_local_degrees( - std::tuple vertex_partition_range, rmm::cuda_stream_view stream) const + std::tuple vertex_partition_range, cuda::stream_ref stream) const { return detail::compute_local_degrees_with_mask_sg( cuda::std::nullopt, vertex_partition_range, this->offsets_, stream); @@ -1137,7 +1137,7 @@ class edge_partition_device_view_t, int> = 0> __host__ rmm::device_uvector compute_local_degrees(MajorIterator major_first, MajorIterator major_last, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { return detail::compute_local_degrees_with_mask_sg( cuda::std::nullopt, major_first, major_last, this->offsets_, stream); diff --git a/cpp/include/cugraph/edge_src_dst_property.hpp b/cpp/include/cugraph/edge_src_dst_property.hpp index 55ff599a52..e36c6adca5 100644 --- a/cpp/include/cugraph/edge_src_dst_property.hpp +++ b/cpp/include/cugraph/edge_src_dst_property.hpp @@ -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 */ @@ -17,6 +17,7 @@ #include #include +#include #include #include @@ -418,7 +419,7 @@ class edge_minor_property_t { key_chunk_start_offsets_.reset(); key_chunk_size_.reset(); - rmm::cuda_stream_view stream{}; + cuda::stream_ref stream{}; if constexpr (std::is_arithmetic_v) { stream = buffer_.stream(); } else { diff --git a/cpp/include/cugraph/host_staging_buffer_manager.hpp b/cpp/include/cugraph/host_staging_buffer_manager.hpp index f906a6bd43..03f4ff71b3 100644 --- a/cpp/include/cugraph/host_staging_buffer_manager.hpp +++ b/cpp/include/cugraph/host_staging_buffer_manager.hpp @@ -1,5 +1,5 @@ /* - * 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 */ #pragma once @@ -9,11 +9,12 @@ #include -#include #include #include #include +#include + #include #include @@ -47,7 +48,7 @@ class host_staging_buffer_manager { } template - static rmm::device_uvector allocate_staging_buffer(size_t size, rmm::cuda_stream_view stream) + static rmm::device_uvector allocate_staging_buffer(size_t size, cuda::stream_ref stream) { auto& s = state(); return rmm::device_uvector(size, stream, *s.pinned_pool_mr); diff --git a/cpp/include/cugraph/large_buffer_manager.hpp b/cpp/include/cugraph/large_buffer_manager.hpp index 1d4fcc2ed2..6d53639d1f 100644 --- a/cpp/include/cugraph/large_buffer_manager.hpp +++ b/cpp/include/cugraph/large_buffer_manager.hpp @@ -1,5 +1,5 @@ /* - * 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 */ #pragma once @@ -10,11 +10,12 @@ #include -#include #include #include #include +#include + #include namespace CUGRAPH_EXPORT cugraph { @@ -52,15 +53,14 @@ enum class large_buffer_type_t { MEMORY, STORAGE, NUM_TYPES }; class large_buffer_manager { public: template - static dataframe_buffer_type_t allocate_memory_buffer(size_t size, - rmm::cuda_stream_view stream) + static dataframe_buffer_type_t allocate_memory_buffer(size_t size, cuda::stream_ref stream) { CUGRAPH_EXPECTS(memory_buffer_initialized(), "large memory buffer resource is not set."); return allocate_dataframe_buffer(size, stream, memory_buffer_resource()->get()); } template - static storage_buffer_t allocate_storage_buffer(size_t, rmm::cuda_stream_view) + static storage_buffer_t allocate_storage_buffer(size_t, cuda::stream_ref) { CUGRAPH_EXPECTS(storage_buffer_initialized(), "large storage buffer resource is not set."); return storage_buffer_t(); diff --git a/cpp/include/cugraph/legacy/functions.hpp b/cpp/include/cugraph/legacy/functions.hpp index 2ccb676c23..141a2f5731 100644 --- a/cpp/include/cugraph/legacy/functions.hpp +++ b/cpp/include/cugraph/legacy/functions.hpp @@ -64,7 +64,7 @@ std::unique_ptr> coo_to_csr( template void comms_bcast(const raft::handle_t& handle, value_t* value, size_t count) { - handle.get_comms().bcast(value, count, 0, handle.get_stream()); + handle.get_comms().bcast(value, count, 0, handle.get_stream().get()); } } // namespace CUGRAPH_EXPORT cugraph diff --git a/cpp/include/cugraph/mtmg/detail/per_device_edgelist.hpp b/cpp/include/cugraph/mtmg/detail/per_device_edgelist.hpp index 32b233157c..ddbd7e87de 100644 --- a/cpp/include/cugraph/mtmg/detail/per_device_edgelist.hpp +++ b/cpp/include/cugraph/mtmg/detail/per_device_edgelist.hpp @@ -15,6 +15,8 @@ #include +#include + namespace CUGRAPH_EXPORT cugraph { namespace mtmg { @@ -78,7 +80,7 @@ class per_device_edgelist_t { */ per_device_edgelist_t(size_t device_buffer_size, std::vector const& edge_property_types, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) : device_buffer_size_{device_buffer_size}, current_pos_{0}, src_{}, @@ -119,7 +121,7 @@ class per_device_edgelist_t { void append(raft::host_span src, raft::host_span dst, raft::host_span edge_properties, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { CUGRAPH_EXPECTS(edge_properties.size() == edge_property_buffers_.size(), "Edge property count mismatch"); @@ -198,7 +200,7 @@ class per_device_edgelist_t { * * @param stream_view CUDA stream view */ - void finalize_buffer(rmm::cuda_stream_view stream_view) + void finalize_buffer(cuda::stream_ref stream_view) { src_.back().resize(current_pos_, stream_view); dst_.back().resize(current_pos_, stream_view); @@ -295,7 +297,7 @@ class per_device_edgelist_t { std::vector resize_and_copy_buffers( std::vector&& buffer, size_t total_size, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { return cugraph::variant_type_dispatch(buffer[0], [&buffer, total_size, stream](auto& buffer0) { using T = typename std::decay_t::value_type; @@ -319,7 +321,7 @@ class per_device_edgelist_t { template std::vector> resize_and_copy_buffers( - std::vector>&& buffer, size_t total_size, rmm::cuda_stream_view stream) + std::vector>&& buffer, size_t total_size, cuda::stream_ref stream) { rmm::device_uvector new_buffer(total_size, stream); @@ -336,7 +338,7 @@ class per_device_edgelist_t { return result; } - void create_new_buffers(rmm::cuda_stream_view stream_view) + void create_new_buffers(cuda::stream_ref stream_view) { src_.emplace_back(device_buffer_size_, stream_view); dst_.emplace_back(device_buffer_size_, stream_view); diff --git a/cpp/include/cugraph/mtmg/handle.hpp b/cpp/include/cugraph/mtmg/handle.hpp index cb6faf0d15..fc124dd2d9 100644 --- a/cpp/include/cugraph/mtmg/handle.hpp +++ b/cpp/include/cugraph/mtmg/handle.hpp @@ -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,6 +11,8 @@ #include +#include + namespace CUGRAPH_EXPORT cugraph { namespace mtmg { @@ -51,7 +53,7 @@ class handle_t { * * @return cuda stream */ - rmm::cuda_stream_view get_stream() const + cuda::stream_ref get_stream() const { return raft_handle_.is_stream_pool_initialized() ? raft_handle_.get_stream_from_stream_pool(thread_rank_) @@ -63,7 +65,7 @@ class handle_t { * * @param stream Which stream to synchronize (defaults to the stream for this handle) */ - void sync_stream(rmm::cuda_stream_view stream) const { raft_handle_.sync_stream(stream); } + void sync_stream(cuda::stream_ref stream) const { raft_handle_.sync_stream(stream); } /** * @brief Sync on the cuda stream for this handle @@ -82,7 +84,7 @@ class handle_t { * * @return exec policy using the current stream */ - rmm::exec_policy get_thrust_policy(rmm::cuda_stream_view stream) const + rmm::exec_policy get_thrust_policy(cuda::stream_ref stream) const { return rmm::exec_policy(stream); } diff --git a/cpp/include/cugraph/mtmg/per_thread_edgelist.hpp b/cpp/include/cugraph/mtmg/per_thread_edgelist.hpp index b1be2cba59..b214dd682a 100644 --- a/cpp/include/cugraph/mtmg/per_thread_edgelist.hpp +++ b/cpp/include/cugraph/mtmg/per_thread_edgelist.hpp @@ -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 */ @@ -9,6 +9,8 @@ #include #include +#include + namespace CUGRAPH_EXPORT cugraph { namespace mtmg { @@ -66,7 +68,7 @@ class per_thread_edgelist_t { void append(vertex_t src, vertex_t dst, std::vector edge_properties, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { if (current_pos_ == src_.size()) { flush(stream_view); } @@ -97,7 +99,7 @@ class per_thread_edgelist_t { void append(raft::host_span src, raft::host_span dst, std::vector> edge_properties, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { size_t count = src.size(); size_t pos = 0; @@ -134,7 +136,7 @@ class per_thread_edgelist_t { * @param sync If true, synchronize the asynchronous copy of data; * defaults to false. */ - void flush(rmm::cuda_stream_view stream_view, bool sync = false) + void flush(cuda::stream_ref stream_view, bool sync = false) { std::vector edge_properties_spans; std::for_each(edge_properties_.begin(), @@ -155,7 +157,7 @@ class per_thread_edgelist_t { current_pos_ = 0; - if (sync) stream_view.synchronize(); + if (sync) stream_view.sync(); } private: diff --git a/cpp/include/cugraph/mtmg/resource_manager.hpp b/cpp/include/cugraph/mtmg/resource_manager.hpp index c2c7ca1113..d612929698 100644 --- a/cpp/include/cugraph/mtmg/resource_manager.hpp +++ b/cpp/include/cugraph/mtmg/resource_manager.hpp @@ -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 */ @@ -17,6 +17,7 @@ #include #include +#include #include @@ -182,7 +183,7 @@ class resource_manager_t { nccl_comms.push_back(std::make_unique()); handles.push_back( - std::make_unique(rmm::cuda_stream_per_thread, + std::make_unique(cuda::stream_ref{cudaStreamPerThread}, std::make_shared(n_streams), per_device_rmm_resources_.find(rank)->second)); device_ids.push_back(pos->second); diff --git a/cpp/include/cugraph/partition_manager.hpp b/cpp/include/cugraph/partition_manager.hpp index 1c6411bfb1..97d79dd224 100644 --- a/cpp/include/cugraph/partition_manager.hpp +++ b/cpp/include/cugraph/partition_manager.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/include/cugraph/prims/count_if_v.cuh b/cpp/include/cugraph/prims/count_if_v.cuh index be056239e9..476d984f18 100644 --- a/cpp/include/cugraph/prims/count_if_v.cuh +++ b/cpp/include/cugraph/prims/count_if_v.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/include/cugraph/prims/detail/extract_transform_if_v_frontier_e.cuh b/cpp/include/cugraph/prims/detail/extract_transform_if_v_frontier_e.cuh index 639305b47f..3fcc667cfa 100644 --- a/cpp/include/cugraph/prims/detail/extract_transform_if_v_frontier_e.cuh +++ b/cpp/include/cugraph/prims/detail/extract_transform_if_v_frontier_e.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -38,6 +38,7 @@ #include #include #include +#include #include #include #include @@ -563,7 +564,7 @@ void extract_transform_if_v_frontier_e_edge_partition( extract_transform_if_v_frontier_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); extract_transform_if_v_frontier_e_high_degree - <<>>( + <<>>( edge_partition, edge_partition_frontier_key_first, raft::device_span((*high_segment_key_local_degree_offsets).data(), @@ -587,7 +588,7 @@ void extract_transform_if_v_frontier_e_edge_partition( extract_transform_if_v_frontier_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); extract_transform_if_v_frontier_e_mid_degree - <<>>( + <<>>( edge_partition, edge_partition_frontier_key_first + (*key_segment_offsets)[1], edge_partition_frontier_key_first + (*key_segment_offsets)[2], @@ -610,7 +611,7 @@ void extract_transform_if_v_frontier_e_edge_partition( extract_transform_if_v_frontier_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); extract_transform_if_v_frontier_e_hypersparse_or_low_degree - <<>>( + <<>>( edge_partition, edge_partition_frontier_key_first + (*key_segment_offsets)[2], edge_partition_frontier_key_first + (*key_segment_offsets)[3], @@ -634,7 +635,7 @@ void extract_transform_if_v_frontier_e_edge_partition( extract_transform_if_v_frontier_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); extract_transform_if_v_frontier_e_hypersparse_or_low_degree - <<>>( + <<>>( edge_partition, edge_partition_frontier_key_first + (*key_segment_offsets)[3], edge_partition_frontier_key_first + (*key_segment_offsets)[4], @@ -662,7 +663,7 @@ void extract_transform_if_v_frontier_e_edge_partition( handle.get_device_properties().maxGridSize[0]); extract_transform_if_v_frontier_e_hypersparse_or_low_degree - <<>>( + <<>>( edge_partition, edge_partition_frontier_key_first, edge_partition_frontier_key_last, diff --git a/cpp/include/cugraph/prims/detail/multi_stream_utils.cuh b/cpp/include/cugraph/prims/detail/multi_stream_utils.cuh index a5cde2a12b..b16d9b489a 100644 --- a/cpp/include/cugraph/prims/detail/multi_stream_utils.cuh +++ b/cpp/include/cugraph/prims/detail/multi_stream_utils.cuh @@ -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 */ #pragma once @@ -10,12 +10,12 @@ #include #include -#include #include #include #include #include +#include #include #include @@ -54,7 +54,7 @@ void copy_if_nosync(InputIterator input_first, FlagIterator flag_first, OutputIterator output_first, raft::device_span count /* size = 1 */, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { CUGRAPH_EXPECTS( static_cast(cuda::std::distance(input_first, input_last)) <= @@ -73,7 +73,7 @@ void copy_if_nosync(InputIterator input_first, output_first, count.data(), input_size, - stream_view); + stream_view.get()); auto d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, stream_view); @@ -84,7 +84,7 @@ void copy_if_nosync(InputIterator input_first, output_first, count.data(), input_size, - stream_view); + stream_view.get()); } template @@ -92,7 +92,7 @@ void count_nosync(InputIterator input_first, InputIterator input_last, raft::device_span count /* size = 1 */, typename thrust::iterator_traits::value_type value, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { CUGRAPH_EXPECTS( static_cast(cuda::std::distance(input_first, input_last)) <= @@ -121,7 +121,7 @@ void sum_nosync( InputIterator input_first, InputIterator input_last, raft::device_span::value_type> sum /* size = 1 */, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { CUGRAPH_EXPECTS( static_cast(cuda::std::distance(input_first, input_last)) <= diff --git a/cpp/include/cugraph/prims/detail/nbr_intersection.cuh b/cpp/include/cugraph/prims/detail/nbr_intersection.cuh index a64f59ac2e..7c22626979 100644 --- a/cpp/include/cugraph/prims/detail/nbr_intersection.cuh +++ b/cpp/include/cugraph/prims/detail/nbr_intersection.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/include/cugraph/prims/detail/optional_dataframe_buffer.hpp b/cpp/include/cugraph/prims/detail/optional_dataframe_buffer.hpp index 5a5d6a0da2..65199330c4 100644 --- a/cpp/include/cugraph/prims/detail/optional_dataframe_buffer.hpp +++ b/cpp/include/cugraph/prims/detail/optional_dataframe_buffer.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -7,6 +7,8 @@ #include #include +#include + #include namespace CUGRAPH_EXPORT cugraph { @@ -33,7 +35,7 @@ struct optional_dataframe_buffer_iterator_value_type_t< }; template -auto allocate_optional_dataframe_buffer(size_t size, rmm::cuda_stream_view stream) +auto allocate_optional_dataframe_buffer(size_t size, cuda::stream_ref stream) { if constexpr (std::is_same_v) { return std::byte{0}; // dummy @@ -44,7 +46,7 @@ auto allocate_optional_dataframe_buffer(size_t size, rmm::cuda_stream_view strea template struct optional_dataframe_buffer_type { - using type = decltype(allocate_optional_dataframe_buffer(size_t{0}, rmm::cuda_stream_view{})); + using type = decltype(allocate_optional_dataframe_buffer(size_t{0}, cuda::stream_ref{})); }; template @@ -98,7 +100,7 @@ template void reserve_optional_dataframe_buffer( optional_dataframe_buffer_type_t& optional_dataframe_buffer, size_t new_buffer_capacity, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { if constexpr (std::is_same_v) { return; @@ -111,7 +113,7 @@ template void resize_optional_dataframe_buffer( optional_dataframe_buffer_type_t& optional_dataframe_buffer, size_t new_buffer_size, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { if constexpr (std::is_same_v) { return; @@ -122,7 +124,7 @@ void resize_optional_dataframe_buffer( template void shrink_to_fit_optional_dataframe_buffer( - optional_dataframe_buffer_type_t& optional_dataframe_buffer, rmm::cuda_stream_view stream_view) + optional_dataframe_buffer_type_t& optional_dataframe_buffer, cuda::stream_ref stream_view) { if constexpr (std::is_same_v) { return; diff --git a/cpp/include/cugraph/prims/detail/per_v_transform_reduce_e.cuh b/cpp/include/cugraph/prims/detail/per_v_transform_reduce_e.cuh index 2f1f97b9b8..2ec53c87c2 100644 --- a/cpp/include/cugraph/prims/detail/per_v_transform_reduce_e.cuh +++ b/cpp/include/cugraph/prims/detail/per_v_transform_reduce_e.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -44,6 +44,7 @@ #include #include #include +#include #include #include #include @@ -895,7 +896,7 @@ void copy_valid_offset_value_pairs( raft::device_span::value_type const>> hypersparse_key_offsets, typename thrust::iterator_traits::value_type invalid_value, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { using offset_t = std::decay_t::value_type>; using value_t = std::decay_t::value_type>; @@ -1033,7 +1034,7 @@ void per_v_transform_reduce_e_edge_partition( assert(segment_key_last == nullptr); } detail::per_v_transform_reduce_e_hypersparse - <<>>( + <<>>( edge_partition, segment_key_first, segment_key_last, @@ -1068,7 +1069,7 @@ void per_v_transform_reduce_e_edge_partition( } *segment_key_first += (*key_segment_offsets)[2]; detail::per_v_transform_reduce_e_low_degree - <<>>( + <<>>( edge_partition, *segment_key_first, *segment_key_first + ((*key_segment_offsets)[3] - (*key_segment_offsets)[2]), @@ -1102,7 +1103,7 @@ void per_v_transform_reduce_e_edge_partition( } *segment_key_first += (*key_segment_offsets)[1]; detail::per_v_transform_reduce_e_mid_degree - <<>>( + <<>>( edge_partition, *segment_key_first, *segment_key_first + ((*key_segment_offsets)[2] - (*key_segment_offsets)[1]), @@ -1137,7 +1138,7 @@ void per_v_transform_reduce_e_edge_partition( segment_key_first = thrust::make_counting_iterator(edge_partition.major_range_first()); } detail::per_v_transform_reduce_e_high_degree - <<>>( + <<>>( edge_partition, *segment_key_first, *segment_key_first + (*key_segment_offsets)[1], @@ -1179,7 +1180,7 @@ void per_v_transform_reduce_e_edge_partition( segment_key_first = thrust::make_counting_iterator(edge_partition.major_range_first()); } detail::per_v_transform_reduce_e_low_degree - <<>>( + <<>>( edge_partition, *segment_key_first, *segment_key_first + num_keys, diff --git a/cpp/include/cugraph/prims/detail/sample_and_compute_local_nbr_indices.cuh b/cpp/include/cugraph/prims/detail/sample_and_compute_local_nbr_indices.cuh index 769bc65a61..88c27ebb43 100644 --- a/cpp/include/cugraph/prims/detail/sample_and_compute_local_nbr_indices.cuh +++ b/cpp/include/cugraph/prims/detail/sample_and_compute_local_nbr_indices.cuh @@ -612,7 +612,7 @@ compute_valid_local_nbr_count_inclusive_sums(raft::handle_t const& handle, compute_valid_local_nbr_count_inclusive_sums_mid_local_degree<<>>( + handle.get_stream().get()>>>( edge_partition, *edge_partition_e_mask, aggregate_local_frontier_major_first + local_frontier_offsets[i], @@ -631,7 +631,7 @@ compute_valid_local_nbr_count_inclusive_sums(raft::handle_t const& handle, compute_valid_local_nbr_count_inclusive_sums_high_local_degree<<>>( + handle.get_stream().get()>>>( edge_partition, *edge_partition_e_mask, aggregate_local_frontier_major_first + local_frontier_offsets[i], diff --git a/cpp/include/cugraph/prims/detail/transform_v_frontier_e.cuh b/cpp/include/cugraph/prims/detail/transform_v_frontier_e.cuh index ca1d97b048..a97ddcaf79 100644 --- a/cpp/include/cugraph/prims/detail/transform_v_frontier_e.cuh +++ b/cpp/include/cugraph/prims/detail/transform_v_frontier_e.cuh @@ -540,7 +540,7 @@ auto transform_v_frontier_e(raft::handle_t const& handle, detail::transform_v_frontier_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_v_frontier_e_high_degree - <<>>( + <<>>( edge_partition, edge_partition_frontier_key_first, edge_partition_key_indices.begin() + edge_partition_v_frontier_partition_offsets[0], @@ -560,7 +560,7 @@ auto transform_v_frontier_e(raft::handle_t const& handle, detail::transform_v_frontier_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_v_frontier_e_mid_degree - <<>>( + <<>>( edge_partition, edge_partition_frontier_key_first, edge_partition_key_indices.begin() + edge_partition_v_frontier_partition_offsets[1], @@ -580,7 +580,7 @@ auto transform_v_frontier_e(raft::handle_t const& handle, detail::transform_v_frontier_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_v_frontier_e_hypersparse_or_low_degree - <<>>( + <<>>( edge_partition, edge_partition_frontier_key_first, edge_partition_key_indices.begin() + edge_partition_v_frontier_partition_offsets[2], @@ -600,7 +600,7 @@ auto transform_v_frontier_e(raft::handle_t const& handle, detail::transform_v_frontier_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_v_frontier_e_hypersparse_or_low_degree - <<>>( + <<>>( edge_partition, edge_partition_frontier_key_first, edge_partition_key_indices.begin() + edge_partition_v_frontier_partition_offsets[3], @@ -621,7 +621,7 @@ auto transform_v_frontier_e(raft::handle_t const& handle, handle.get_device_properties().maxGridSize[0]); detail::transform_v_frontier_e_hypersparse_or_low_degree - <<>>( + <<>>( edge_partition, edge_partition_frontier_key_first, thrust::make_counting_iterator(size_t{0}), diff --git a/cpp/include/cugraph/prims/fill_edge_src_dst_property.cuh b/cpp/include/cugraph/prims/fill_edge_src_dst_property.cuh index 39bea2bddc..e892b375af 100644 --- a/cpp/include/cugraph/prims/fill_edge_src_dst_property.cuh +++ b/cpp/include/cugraph/prims/fill_edge_src_dst_property.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -27,6 +27,7 @@ #include #include #include +#include #include #include diff --git a/cpp/include/cugraph/prims/key_store.cuh b/cpp/include/cugraph/prims/key_store.cuh index ec53e62dc2..c0b0fcb557 100644 --- a/cpp/include/cugraph/prims/key_store.cuh +++ b/cpp/include/cugraph/prims/key_store.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 */ #pragma once @@ -13,6 +13,7 @@ #include #include #include +#include #include #include #include @@ -116,7 +117,7 @@ class key_binary_search_store_view_t { void contains(QueryKeyIterator key_first, QueryKeyIterator key_last, ResultValueIterator value_first, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { thrust::transform( rmm::exec_policy(stream), @@ -157,7 +158,7 @@ class key_cuco_store_view_t { void contains(QueryKeyIterator key_first, QueryKeyIterator key_last, ResultValueIterator value_first, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { cuco_store_->contains(key_first, key_last, value_first, stream); } @@ -177,7 +178,7 @@ class key_binary_search_store_t { public: using key_type = key_t; - key_binary_search_store_t(rmm::cuda_stream_view stream) : store_keys_(0, stream) {} + key_binary_search_store_t(cuda::stream_ref stream) : store_keys_(0, stream) {} template key_binary_search_store_t( @@ -186,7 +187,7 @@ class key_binary_search_store_t { bool key_sorted /* if set to true, assume that the input data is sorted and skip sorting (which is necessary for binary-search) */ , - rmm::cuda_stream_view stream) + cuda::stream_ref stream) : store_keys_(static_cast(cuda::std::distance(key_first, key_last)), stream) { thrust::copy(rmm::exec_policy(stream), key_first, key_last, store_keys_.begin()); @@ -200,7 +201,7 @@ class key_binary_search_store_t { bool key_sorted /* if set to true, assume that the input data is sorted and skip sorting (which is necessary for binary-search) */ , - rmm::cuda_stream_view stream) + cuda::stream_ref stream) : store_keys_(std::move(keys)) { if (!key_sorted) { @@ -208,7 +209,7 @@ class key_binary_search_store_t { } } - auto release(rmm::cuda_stream_view stream) + auto release(cuda::stream_ref stream) { auto tmp_store_keys = std::move(store_keys_); store_keys_ = rmm::device_uvector(0, stream); @@ -241,9 +242,9 @@ class key_cuco_store_t { rmm::mr::polymorphic_allocator, cuco_storage_type>; - key_cuco_store_t(rmm::cuda_stream_view stream) {} + key_cuco_store_t(cuda::stream_ref stream) {} - key_cuco_store_t(size_t capacity, key_t invalid_key, rmm::cuda_stream_view stream) + key_cuco_store_t(size_t capacity, key_t invalid_key, cuda::stream_ref stream) { allocate(capacity, invalid_key, stream); capacity_ = capacity; @@ -254,7 +255,7 @@ class key_cuco_store_t { key_cuco_store_t(KeyIterator key_first, KeyIterator key_last, key_t invalid_key, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto num_keys = static_cast(cuda::std::distance(key_first, key_last)); allocate(num_keys, invalid_key, stream); @@ -265,12 +266,12 @@ class key_cuco_store_t { } template - void insert(KeyIterator key_first, KeyIterator key_last, rmm::cuda_stream_view stream) + void insert(KeyIterator key_first, KeyIterator key_last, cuda::stream_ref stream) { auto num_keys = static_cast(cuda::std::distance(key_first, key_last)); if (num_keys == 0) return; - size_ += cuco_store_->insert(key_first, key_last, stream.value()); + size_ += cuco_store_->insert(key_first, key_last, stream.get()); } template @@ -278,18 +279,18 @@ class key_cuco_store_t { KeyIterator key_last, StencilIterator stencil_first, PredOp pred_op, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto num_keys = static_cast(cuda::std::distance(key_first, key_last)); if (num_keys == 0) return; - size_ += cuco_store_->insert_if(key_first, key_last, stencil_first, pred_op, stream.value()); + size_ += cuco_store_->insert_if(key_first, key_last, stencil_first, pred_op, stream.get()); } - auto release(rmm::cuda_stream_view stream) + auto release(cuda::stream_ref stream) { rmm::device_uvector keys(size(), stream); - auto last = cuco_store_->retrieve_all(keys.begin(), stream.value()); + auto last = cuco_store_->retrieve_all(keys.begin(), stream.get()); keys.resize(cuda::std::distance(keys.begin(), last), stream); keys.shrink_to_fit(stream); allocate(0, invalid_key(), stream); @@ -307,7 +308,7 @@ class key_cuco_store_t { size_t capacity() const { return capacity_; } private: - void allocate(size_t num_keys, key_t invalid_key, rmm::cuda_stream_view stream) + void allocate(size_t num_keys, key_t invalid_key, cuda::stream_ref stream) { double constexpr load_factor = 0.7; auto cuco_size = std::max( @@ -323,7 +324,7 @@ class key_cuco_store_t { cuco::thread_scope_device, cuco_storage_type{}, rmm::mr::polymorphic_allocator{}, - stream.value()); + stream.get()); } std::unique_ptr cuco_store_{nullptr}; @@ -344,7 +345,7 @@ class key_store_t { static_assert(std::is_arithmetic_v); - key_store_t(rmm::cuda_stream_view stream) : store_(stream) {} + key_store_t(cuda::stream_ref stream) : store_(stream) {} /* when use_binary_search = false */ template @@ -353,7 +354,7 @@ class key_store_t { capacity can be larger (for performance & correctness reasons) */ , key_t invalid_key /* invalid key shouldn't appear in any *iter in [key_first, key_last) */, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, std::enable_if_t = 0) : store_(capacity, invalid_key, stream) { @@ -366,7 +367,7 @@ class key_store_t { bool key_sorted /* if set to true, assume that the input data is sorted and skip sorting (which is necessary for binary-search) */ , - rmm::cuda_stream_view stream, + cuda::stream_ref stream, std::enable_if_t = 0) : store_(key_first, key_last, key_sorted, stream) { @@ -378,7 +379,7 @@ class key_store_t { KeyIterator key_first, KeyIterator key_last, key_t invalid_key /* invalid key shouldn't appear in any *iter in [key_first, key_last) */, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, std::enable_if_t = 0) : store_(key_first, key_last, invalid_key, stream) { @@ -390,7 +391,7 @@ class key_store_t { bool key_sorted /* if set to true, assume that the input data is sorted and skip sorting (which is necessary for binary-search) */ , - rmm::cuda_stream_view stream, + cuda::stream_ref stream, std::enable_if_t = 0) : store_(std::move(keys), key_sorted, stream) { @@ -400,7 +401,7 @@ class key_store_t { template std::enable_if_t insert(KeyIterator key_first, KeyIterator key_last, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { store_.insert(key_first, key_last, stream); } @@ -414,13 +415,13 @@ class key_store_t { KeyIterator key_last, StencilIterator stencil_first, PredOp pred_op, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { store_.insert_if(key_first, key_last, stencil_first, pred_op, stream); } // key_store_t becomes empty after release - auto release(rmm::cuda_stream_view stream) { return store_.release(stream); } + auto release(cuda::stream_ref stream) { return store_.release(stream); } auto view() const { diff --git a/cpp/include/cugraph/prims/kv_store.cuh b/cpp/include/cugraph/prims/kv_store.cuh index 49f457fb5e..28bc641c30 100644 --- a/cpp/include/cugraph/prims/kv_store.cuh +++ b/cpp/include/cugraph/prims/kv_store.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -23,6 +23,7 @@ #include #include #include +#include #include #include #include @@ -253,7 +254,7 @@ class kv_binary_search_store_view_t { void find(QueryKeyIterator key_first, QueryKeyIterator key_last, ResultValueIterator value_first, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { thrust::transform(rmm::exec_policy(stream), key_first, @@ -267,7 +268,7 @@ class kv_binary_search_store_view_t { void contains(QueryKeyIterator key_first, QueryKeyIterator key_last, ResultValueIterator value_first, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { thrust::transform( rmm::exec_policy(stream), @@ -333,14 +334,14 @@ class kv_cuco_store_view_t { void find(QueryKeyIterator key_first, QueryKeyIterator key_last, ResultValueIterator value_first, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { if constexpr (std::is_arithmetic_v) { - cuco_store_->find(key_first, key_last, value_first, stream.value()); + cuco_store_->find(key_first, key_last, value_first, stream.get()); } else { rmm::device_uvector indices(cuda::std::distance(key_first, key_last), stream); auto invalid_idx = cuco_store_->empty_value_sentinel(); - cuco_store_->find(key_first, key_last, indices.begin(), stream.value()); + cuco_store_->find(key_first, key_last, indices.begin(), stream.get()); thrust::transform(rmm::exec_policy(stream), indices.begin(), indices.end(), @@ -354,9 +355,9 @@ class kv_cuco_store_view_t { void contains(QueryKeyIterator key_first, QueryKeyIterator key_last, ResultValueIterator value_first, - rmm::cuda_stream_view stream) const + cuda::stream_ref stream) const { - cuco_store_->contains(key_first, key_last, value_first, stream.value()); + cuco_store_->contains(key_first, key_last, value_first, stream.get()); } auto cuco_store_find_device_ref() const { return cuco_store_->ref(cuco::find); } @@ -393,7 +394,7 @@ class kv_binary_search_store_t { using key_type = key_t; using value_type = value_t; - kv_binary_search_store_t(rmm::cuda_stream_view stream) + kv_binary_search_store_t(cuda::stream_ref stream) : store_keys_(0, stream), store_values_(allocate_dataframe_buffer(0, stream)) { } @@ -407,7 +408,7 @@ class kv_binary_search_store_t { bool key_sorted /* if set to true, assume that the input data is sorted and skip sorting (which is necessary for binary-search) */ , - rmm::cuda_stream_view stream) + cuda::stream_ref stream) : store_keys_(static_cast(cuda::std::distance(key_first, key_last)), stream), store_values_(allocate_dataframe_buffer( static_cast(cuda::std::distance(key_first, key_last)), stream)), @@ -429,12 +430,12 @@ class kv_binary_search_store_t { kv_binary_search_store_t( rmm::device_uvector&& keys, - decltype(allocate_dataframe_buffer(0, rmm::cuda_stream_view{}))&& values, + decltype(allocate_dataframe_buffer(0, cuda::stream_ref{}))&& values, value_t invalid_value /* invalid_value is returned when match fails for the given key */, bool key_sorted /* if set to true, assume that the input data is sorted and skip sorting (which is necessary for binary-search) */ , - rmm::cuda_stream_view stream) + cuda::stream_ref stream) : store_keys_(std::move(keys)), store_values_(std::move(values)), invalid_value_(invalid_value) { if (!key_sorted) { @@ -445,7 +446,7 @@ class kv_binary_search_store_t { } } - auto retrieve_all(rmm::cuda_stream_view stream) + auto retrieve_all(cuda::stream_ref stream) { rmm::device_uvector tmp_store_keys(store_keys_.size(), stream); auto tmp_store_values = @@ -459,7 +460,7 @@ class kv_binary_search_store_t { return std::make_tuple(std::move(tmp_store_keys), std::move(tmp_store_values)); } - auto release(rmm::cuda_stream_view stream) + auto release(cuda::stream_ref stream) { auto tmp_store_keys = std::move(store_keys_); auto tmp_store_values = std::move(store_values_); @@ -482,7 +483,7 @@ class kv_binary_search_store_t { private: rmm::device_uvector store_keys_; - decltype(allocate_dataframe_buffer(0, rmm::cuda_stream_view{})) store_values_; + decltype(allocate_dataframe_buffer(0, cuda::stream_ref{})) store_values_; value_t invalid_value_{}; }; @@ -490,10 +491,9 @@ class kv_binary_search_store_t { template class kv_cuco_store_t { public: - using key_type = key_t; - using value_type = value_t; - using value_buffer_type = - decltype(allocate_dataframe_buffer(0, rmm::cuda_stream_view{})); + using key_type = key_t; + using value_type = value_t; + using value_buffer_type = decltype(allocate_dataframe_buffer(0, cuda::stream_ref{})); using const_value_iterator = std::invoke_result_t), value_buffer_type&>; @@ -509,7 +509,7 @@ class kv_cuco_store_t { rmm::mr::polymorphic_allocator, cuco_storage_type>; - kv_cuco_store_t(rmm::cuda_stream_view stream) + kv_cuco_store_t(cuda::stream_ref stream) : store_values_(allocate_optional_dataframe_buffer< std::conditional_t, value_t, void>>(0, stream)) { @@ -518,7 +518,7 @@ class kv_cuco_store_t { kv_cuco_store_t(size_t capacity, key_t invalid_key, value_t invalid_value, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) : store_values_(allocate_optional_dataframe_buffer< std::conditional_t, value_t, void>>(0, stream)) { @@ -534,7 +534,7 @@ class kv_cuco_store_t { ValueIterator value_first, key_t invalid_key, value_t invalid_value, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) : store_values_(allocate_optional_dataframe_buffer< std::conditional_t, value_t, void>>(0, stream)) { @@ -555,7 +555,7 @@ class kv_cuco_store_t { void insert(KeyIterator key_first, KeyIterator key_last, ValueIterator value_first, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { static_assert(std::is_same_v::value_type, key_t>); static_assert( @@ -566,7 +566,7 @@ class kv_cuco_store_t { if constexpr (std::is_arithmetic_v) { auto pair_first = thrust::make_zip_iterator(key_first, value_first); - size_ += cuco_store_->insert(pair_first, pair_first + num_keys, stream.value()); + size_ += cuco_store_->insert(pair_first, pair_first + num_keys, stream.get()); } else { auto old_store_value_size = size_optional_dataframe_buffer(store_values_); // FIXME: we can use cuda::atomic instead but currently on a system with x86 + GPU, this @@ -599,7 +599,7 @@ class kv_cuco_store_t { ValueIterator value_first, StencilIterator stencil_first, PredOp pred_op, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { static_assert(std::is_same_v::value_type, key_t>); static_assert( @@ -611,7 +611,7 @@ class kv_cuco_store_t { if constexpr (std::is_arithmetic_v) { auto pair_first = thrust::make_zip_iterator(key_first, value_first); size_ += cuco_store_->insert_if( - pair_first, pair_first + num_keys, stencil_first, pred_op, stream.value()); + pair_first, pair_first + num_keys, stencil_first, pred_op, stream.get()); } else { auto old_store_value_size = size_optional_dataframe_buffer(store_values_); // FIXME: we can use cuda::atomic instead but currently on a system with x86 + GPU, this @@ -649,7 +649,7 @@ class kv_cuco_store_t { void insert_and_assign(KeyIterator key_first, KeyIterator key_last, ValueIterator value_first, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { static_assert(std::is_same_v::value_type, key_t>); static_assert( @@ -751,7 +751,7 @@ class kv_cuco_store_t { ValueIterator value_first, StencilIterator stencil_first, PredOp pred_op, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto num_keys = static_cast(cuda::std::distance(key_first, key_last)); if (num_keys == 0) return; @@ -792,19 +792,19 @@ class kv_cuco_store_t { } } - auto retrieve_all(rmm::cuda_stream_view stream) + auto retrieve_all(cuda::stream_ref stream) { rmm::device_uvector keys(size_, stream); auto values = allocate_dataframe_buffer(0, stream); if constexpr (std::is_arithmetic_v) { values.resize(size_, stream); - auto pair_last = cuco_store_->retrieve_all(keys.begin(), values.begin(), stream.value()); + auto pair_last = cuco_store_->retrieve_all(keys.begin(), values.begin(), stream.get()); // FIXME: this resize (& shrink_to_fit) shouldn't be necessary if size_ is exact keys.resize(cuda::std::distance(keys.begin(), std::get<0>(pair_last)), stream); values.resize(keys.size(), stream); } else { rmm::device_uvector indices(size_, stream); - auto pair_last = cuco_store_->retrieve_all(keys.begin(), indices.begin(), stream.value()); + auto pair_last = cuco_store_->retrieve_all(keys.begin(), indices.begin(), stream.get()); // FIXME: this resize (& shrink_to_fit) shouldn't be necessary if size_ is exact keys.resize(cuda::std::distance(keys.begin(), std::get<0>(pair_last)), stream); indices.resize(keys.size(), stream); @@ -818,7 +818,7 @@ class kv_cuco_store_t { return std::make_tuple(std::move(keys), std::move(values)); } - auto release(rmm::cuda_stream_view stream) + auto release(cuda::stream_ref stream) { auto [retrieved_keys, retrieved_values] = retrieve_all(stream); allocate(0, invalid_key(), invalid_value(), stream); @@ -852,10 +852,7 @@ class kv_cuco_store_t { size_t capacity() const { return capacity_; } private: - void allocate(size_t num_keys, - key_t invalid_key, - value_t invalid_value, - rmm::cuda_stream_view stream) + void allocate(size_t num_keys, key_t invalid_key, value_t invalid_value, cuda::stream_ref stream) { double constexpr load_factor = 0.7; auto cuco_size = std::max( @@ -873,7 +870,7 @@ class kv_cuco_store_t { cuco::thread_scope_device, cuco_storage_type{}, rmm::mr::polymorphic_allocator{}, - stream.value()); + stream.get()); } else { cuco_store_ = std::make_unique( cuco_size, @@ -885,7 +882,7 @@ class kv_cuco_store_t { cuco::thread_scope_device, cuco_storage_type{}, rmm::mr::polymorphic_allocator{}, - stream.value()); + stream.get()); reserve_optional_dataframe_buffer(store_values_, num_keys, stream); } } @@ -893,7 +890,7 @@ class kv_cuco_store_t { std::unique_ptr cuco_store_{nullptr}; decltype(allocate_optional_dataframe_buffer< std::conditional_t, value_t, void>>( - 0, rmm::cuda_stream_view{})) store_values_; + 0, cuda::stream_ref{})) store_values_; std::conditional_t, value_t, std::byte /* dummy */> invalid_value_{}; @@ -919,7 +916,7 @@ class kv_store_t { static_assert(std::is_arithmetic_v); static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic::value); - kv_store_t(rmm::cuda_stream_view stream) : store_(stream) {} + kv_store_t(cuda::stream_ref stream) : store_(stream) {} /* when use_binary_search = false */ template @@ -932,7 +929,7 @@ class kv_store_t { value_first + cuda::std::distance(key_first, key_last)), invalid_value is returned when match fails for the given key */ , - rmm::cuda_stream_view stream, + cuda::stream_ref stream, std::enable_if_t = 0) : store_(capacity, invalid_key, invalid_value, stream) { @@ -948,7 +945,7 @@ class kv_store_t { bool key_sorted /* if set to true, assume that the input data is sorted and skip sorting (which is necessary for binary-search) */ , - rmm::cuda_stream_view stream, + cuda::stream_ref stream, std::enable_if_t = 0) : store_(key_first, key_last, value_first, invalid_value, key_sorted, stream) { @@ -965,7 +962,7 @@ class kv_store_t { value_first + cuda::std::distance(key_first, key_last)), invalid_value is returned when match fails for the given key */ , - rmm::cuda_stream_view stream, + cuda::stream_ref stream, std::enable_if_t = 0) : store_(key_first, key_last, value_first, invalid_key, invalid_value, stream) { @@ -975,12 +972,12 @@ class kv_store_t { template kv_store_t( rmm::device_uvector&& keys, - decltype(allocate_dataframe_buffer(0, rmm::cuda_stream_view{}))&& values, + decltype(allocate_dataframe_buffer(0, cuda::stream_ref{}))&& values, value_t invalid_value /* invalid_value is returned when match fails for the given key */, bool key_sorted /* if set to true, assume that the input data is sorted and skip sorting (which is necessary for binary-search) */ , - rmm::cuda_stream_view stream, + cuda::stream_ref stream, std::enable_if_t = 0) : store_(std::move(keys), std::move(values), invalid_value, key_sorted, stream) { @@ -991,7 +988,7 @@ class kv_store_t { std::enable_if_t insert(KeyIterator key_first, KeyIterator key_last, ValueIterator value_first, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { store_.insert(key_first, key_last, value_first, stream); } @@ -1007,7 +1004,7 @@ class kv_store_t { ValueIterator value_first, StencilIterator stencil_first, PredOp pred_op, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { store_.insert_if(key_first, key_last, value_first, stencil_first, pred_op, stream); } @@ -1017,7 +1014,7 @@ class kv_store_t { std::enable_if_t insert_and_assign(KeyIterator key_first, KeyIterator key_last, ValueIterator value_first, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { store_.insert_and_assign(key_first, key_last, value_first, stream); } @@ -1033,15 +1030,15 @@ class kv_store_t { ValueIterator value_first, StencilIterator stencil_first, PredOp pred_op, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { store_.insert_and_assign_if(key_first, key_last, value_first, stencil_first, pred_op, stream); } - auto retrieve_all(rmm::cuda_stream_view stream) const { return store_.retrieve_all(stream); } + auto retrieve_all(cuda::stream_ref stream) const { return store_.retrieve_all(stream); } // kv_store_t becomes empty after release - auto release(rmm::cuda_stream_view stream) { return store_.release(stream); } + auto release(cuda::stream_ref stream) { return store_.release(stream); } auto view() const { diff --git a/cpp/include/cugraph/prims/per_v_pair_transform_src_dst_nbr_intersection.cuh b/cpp/include/cugraph/prims/per_v_pair_transform_src_dst_nbr_intersection.cuh index c5e5a374d4..f427a24933 100644 --- a/cpp/include/cugraph/prims/per_v_pair_transform_src_dst_nbr_intersection.cuh +++ b/cpp/include/cugraph/prims/per_v_pair_transform_src_dst_nbr_intersection.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -27,6 +27,7 @@ #include #include #include +#include #include #include #include @@ -206,7 +207,7 @@ void per_v_pair_transform_minor_nbr_intersection( auto num_input_pairs = static_cast(cuda::std::distance(vertex_pair_first, vertex_pair_last)); std::optional> sorted_unique_vertices{std::nullopt}; - std::optional(size_t{0}, rmm::cuda_stream_view{}))> + std::optional(size_t{0}, cuda::stream_ref{}))> property_buffer_for_sorted_unique_vertices{std::nullopt}; if constexpr (GraphViewType::is_multi_gpu) { auto& comm = handle.get_comms(); diff --git a/cpp/include/cugraph/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh b/cpp/include/cugraph/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh index ebc41e2fcf..de30cc1938 100644 --- a/cpp/include/cugraph/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh +++ b/cpp/include/cugraph/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -513,7 +513,7 @@ void per_v_transform_reduce_dst_key_aggregated_outgoing_e( h_vertex_offsets[j + 1] - h_vertex_offsets[j], offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); } else { cub::DeviceSegmentedSort::SortKeys(static_cast(nullptr), tmp_storage_bytes, @@ -544,7 +544,7 @@ void per_v_transform_reduce_dst_key_aggregated_outgoing_e( h_vertex_offsets[j + 1] - h_vertex_offsets[j], offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); } else { cub::DeviceSegmentedSort::SortKeys(d_tmp_storage.data(), tmp_storage_bytes, diff --git a/cpp/include/cugraph/prims/reduce_v.cuh b/cpp/include/cugraph/prims/reduce_v.cuh index a35d91d3be..3643cffc0d 100644 --- a/cpp/include/cugraph/prims/reduce_v.cuh +++ b/cpp/include/cugraph/prims/reduce_v.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/include/cugraph/prims/transform_e.cuh b/cpp/include/cugraph/prims/transform_e.cuh index b544b1d1db..fa9812d79b 100644 --- a/cpp/include/cugraph/prims/transform_e.cuh +++ b/cpp/include/cugraph/prims/transform_e.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 */ #pragma once @@ -353,7 +353,7 @@ void transform_e(raft::handle_t const& handle, handle.get_device_properties().maxGridSize[0]); if (edge_partition_e_mask) { detail::transform_e_packed_bool - <<>>( + <<>>( edge_partition, edge_partition_src_value_input, edge_partition_dst_value_input, @@ -363,7 +363,7 @@ void transform_e(raft::handle_t const& handle, e_op); } else { detail::transform_e_packed_bool - <<>>( + <<>>( edge_partition, edge_partition_src_value_input, edge_partition_dst_value_input, diff --git a/cpp/include/cugraph/prims/transform_reduce_e.cuh b/cpp/include/cugraph/prims/transform_reduce_e.cuh index bebbc544ab..f0e86dde60 100644 --- a/cpp/include/cugraph/prims/transform_reduce_e.cuh +++ b/cpp/include/cugraph/prims/transform_reduce_e.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -507,7 +507,7 @@ T transform_reduce_e(raft::handle_t const& handle, detail::transform_reduce_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_e_high_degree - <<>>( + <<>>( edge_partition, edge_partition.major_range_first(), edge_partition.major_range_first() + (*segment_offsets)[1], @@ -523,7 +523,7 @@ T transform_reduce_e(raft::handle_t const& handle, detail::transform_reduce_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_e_mid_degree - <<>>( + <<>>( edge_partition, edge_partition.major_range_first() + (*segment_offsets)[1], edge_partition.major_range_first() + (*segment_offsets)[2], @@ -539,7 +539,7 @@ T transform_reduce_e(raft::handle_t const& handle, detail::transform_reduce_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_e_low_degree - <<>>( + <<>>( edge_partition, edge_partition.major_range_first() + (*segment_offsets)[2], edge_partition.major_range_first() + (*segment_offsets)[3], @@ -555,7 +555,7 @@ T transform_reduce_e(raft::handle_t const& handle, detail::transform_reduce_e_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_e_hypersparse - <<>>( + <<>>( edge_partition, edge_partition_src_value_input, edge_partition_dst_value_input, @@ -571,7 +571,7 @@ T transform_reduce_e(raft::handle_t const& handle, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_e_low_degree - <<>>( + <<>>( edge_partition, edge_partition.major_range_first(), edge_partition.major_range_last(), diff --git a/cpp/include/cugraph/prims/transform_reduce_e_by_src_dst_key.cuh b/cpp/include/cugraph/prims/transform_reduce_e_by_src_dst_key.cuh index b072f590ae..27a7001b4e 100644 --- a/cpp/include/cugraph/prims/transform_reduce_e_by_src_dst_key.cuh +++ b/cpp/include/cugraph/prims/transform_reduce_e_by_src_dst_key.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -613,7 +613,7 @@ transform_reduce_e_by_src_dst_key(raft::handle_t const& handle, detail::transform_reduce_e_by_src_dst_key_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_by_src_dst_key_high_degree - <<>>( + <<>>( edge_partition, edge_partition.major_range_first(), edge_partition.major_range_first() + (*segment_offsets)[1], @@ -636,7 +636,7 @@ transform_reduce_e_by_src_dst_key(raft::handle_t const& handle, detail::transform_reduce_e_by_src_dst_key_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_by_src_dst_key_mid_degree - <<>>( + <<>>( edge_partition, edge_partition.major_range_first() + (*segment_offsets)[1], edge_partition.major_range_first() + (*segment_offsets)[2], @@ -659,7 +659,7 @@ transform_reduce_e_by_src_dst_key(raft::handle_t const& handle, detail::transform_reduce_e_by_src_dst_key_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_by_src_dst_key_low_degree - <<>>( + <<>>( edge_partition, edge_partition.major_range_first() + (*segment_offsets)[2], edge_partition.major_range_first() + (*segment_offsets)[3], @@ -683,7 +683,7 @@ transform_reduce_e_by_src_dst_key(raft::handle_t const& handle, detail::transform_reduce_e_by_src_dst_key_kernel_block_size, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_by_src_dst_key_hypersparse - <<>>( + <<>>( edge_partition, edge_partition_src_value_input, edge_partition_dst_value_input, @@ -705,7 +705,7 @@ transform_reduce_e_by_src_dst_key(raft::handle_t const& handle, handle.get_device_properties().maxGridSize[0]); detail::transform_reduce_by_src_dst_key_low_degree - <<>>( + <<>>( edge_partition, edge_partition.major_range_first(), edge_partition.major_range_last(), @@ -724,7 +724,7 @@ transform_reduce_e_by_src_dst_key(raft::handle_t const& handle, } } std::tie(tmp_keys, tmp_value_buffer) = reduce_to_unique_kv_pairs( - std::move(tmp_keys), std::move(tmp_value_buffer), reduce_op, handle.get_stream()); + std::move(tmp_keys), std::move(tmp_value_buffer), reduce_op, handle.get_stream().get()); if (GraphViewType::is_multi_gpu) { auto& comm = handle.get_comms(); @@ -753,7 +753,7 @@ transform_reduce_e_by_src_dst_key(raft::handle_t const& handle, reduce_to_unique_kv_pairs(std::move(rx_unique_keys), std::move(rx_value_for_unique_key_buffer), reduce_op, - handle.get_stream()); + handle.get_stream().get()); } auto cur_size = keys.size(); @@ -778,7 +778,7 @@ transform_reduce_e_by_src_dst_key(raft::handle_t const& handle, if (GraphViewType::is_multi_gpu) { std::tie(keys, value_buffer) = reduce_to_unique_kv_pairs( - std::move(keys), std::move(value_buffer), reduce_op, handle.get_stream()); + std::move(keys), std::move(value_buffer), reduce_op, handle.get_stream().get()); } // FIXME: add init diff --git a/cpp/include/cugraph/prims/transform_reduce_src_dst_nbr_intersection_of_e_endpoints_by_v.cuh b/cpp/include/cugraph/prims/transform_reduce_src_dst_nbr_intersection_of_e_endpoints_by_v.cuh index 3c67a6c64a..b842ef78c9 100644 --- a/cpp/include/cugraph/prims/transform_reduce_src_dst_nbr_intersection_of_e_endpoints_by_v.cuh +++ b/cpp/include/cugraph/prims/transform_reduce_src_dst_nbr_intersection_of_e_endpoints_by_v.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/include/cugraph/prims/update_edge_src_dst_property.cuh b/cpp/include/cugraph/prims/update_edge_src_dst_property.cuh index 4b1b337ee9..4b098512c0 100644 --- a/cpp/include/cugraph/prims/update_edge_src_dst_property.cuh +++ b/cpp/include/cugraph/prims/update_edge_src_dst_property.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/include/cugraph/prims/vertex_frontier.cuh b/cpp/include/cugraph/prims/vertex_frontier.cuh index 2f73a9a083..f086dd5dae 100644 --- a/cpp/include/cugraph/prims/vertex_frontier.cuh +++ b/cpp/include/cugraph/prims/vertex_frontier.cuh @@ -27,6 +27,7 @@ #include #include #include +#include #include #include #include @@ -53,7 +54,7 @@ template KeyIterator compute_key_lower_bound(KeyIterator sorted_unique_key_first, KeyIterator sorted_unique_key_last, vertex_t v_threshold, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { using key_t = typename thrust::iterator_traits::value_type; @@ -78,7 +79,7 @@ std::vector compute_key_segment_offsets(KeyIterator sorted_key_first, KeyIterator sorted_key_last, raft::host_span segment_offsets, vertex_t vertex_range_first, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { using key_t = typename thrust::iterator_traits::value_type; @@ -112,7 +113,7 @@ std::vector compute_key_segment_offsets(KeyIterator sorted_key_first, std::vector h_offsets(d_offsets.size() + 2); raft::update_host(h_offsets.data() + 1, d_offsets.data(), d_offsets.size(), stream_view); - RAFT_CUDA_TRY(cudaStreamSynchronize(stream_view)); + stream_view.sync(); h_offsets[0] = size_t{0}; h_offsets.back() = static_cast(cuda::std::distance(sorted_key_first, sorted_key_last)); @@ -125,7 +126,7 @@ rmm::device_uvector compute_vertex_list_bitmap_info( VertexIterator sorted_unique_vertex_last, typename thrust::iterator_traits::value_type vertex_range_first, typename thrust::iterator_traits::value_type vertex_range_last, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { using vertex_t = typename thrust::iterator_traits::value_type; @@ -174,7 +175,7 @@ void device_bcast_vertex_list( typename thrust::iterator_traits::value_type vertex_range_last, size_t v_list_size, int root, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { using vertex_t = typename thrust::iterator_traits::value_type; @@ -214,7 +215,7 @@ void retrieve_vertex_list_from_bitmap( raft::device_span count /* size = 1 */, typename thrust::iterator_traits::value_type vertex_range_first, typename thrust::iterator_traits::value_type vertex_range_last, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { using vertex_t = typename thrust::iterator_traits::value_type; diff --git a/cpp/include/cugraph/utilities/collect_comm.cuh b/cpp/include/cugraph/utilities/collect_comm.cuh index 4e3cfddc8c..d06a33364b 100644 --- a/cpp/include/cugraph/utilities/collect_comm.cuh +++ b/cpp/include/cugraph/utilities/collect_comm.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -24,6 +24,7 @@ #include #include +#include #include #include #include @@ -51,7 +52,7 @@ void find_values_for_collect_keys( ValueIterator collect_value_first, typename KVStoreViewType::key_type invalid_key, typename KVStoreViewType::value_type invalid_value, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { using key_t = typename KVStoreViewType::key_type; using value_t = typename KVStoreViewType::value_type; @@ -109,7 +110,7 @@ collect_values_for_unique_keys( KVStoreViewType kv_store_view, rmm::device_uvector&& collect_unique_keys, KeyToCommRankOp key_to_comm_rank_op, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { using key_t = typename KVStoreViewType::key_type; using value_t = typename KVStoreViewType::value_type; @@ -147,7 +148,7 @@ dataframe_buffer_type_t collect_values_for KeyIterator collect_key_first, KeyIterator collect_key_last, KeyToCommRankOp key_to_comm_rank_op, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { using key_t = typename KVStoreViewType::key_type; static_assert(std::is_same_v::value_type, key_t>); diff --git a/cpp/include/cugraph/utilities/cub_wrappers/device_segmented_sort.cuh b/cpp/include/cugraph/utilities/cub_wrappers/device_segmented_sort.cuh index ee4293e1ff..8a6708f4b9 100644 --- a/cpp/include/cugraph/utilities/cub_wrappers/device_segmented_sort.cuh +++ b/cpp/include/cugraph/utilities/cub_wrappers/device_segmented_sort.cuh @@ -45,7 +45,7 @@ void device_segmented_sort_pairs(raft::handle_t const& handle, begin_offsets.size(), begin_offsets.data(), end_offsets.data(), - handle.get_stream()); + handle.get_stream().get()); rmm::device_uvector d_tmp_storage(tmp_storage_bytes, handle.get_stream()); cub::DeviceSegmentedSort::SortPairs(d_tmp_storage.data(), tmp_storage_bytes, @@ -57,7 +57,7 @@ void device_segmented_sort_pairs(raft::handle_t const& handle, begin_offsets.size(), begin_offsets.data(), end_offsets.data(), - handle.get_stream()); + handle.get_stream().get()); } // offsets is a CSR offset array of size num_segments + 1 diff --git a/cpp/include/cugraph/utilities/dataframe_buffer.hpp b/cpp/include/cugraph/utilities/dataframe_buffer.hpp index 80c1936ce3..c061210d99 100644 --- a/cpp/include/cugraph/utilities/dataframe_buffer.hpp +++ b/cpp/include/cugraph/utilities/dataframe_buffer.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -9,10 +9,10 @@ #include -#include #include #include +#include #include #include @@ -25,7 +25,7 @@ template auto allocate_dataframe_buffer_tuple_impl( std::index_sequence, size_t buffer_size, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()) { return std::make_tuple( @@ -64,7 +64,7 @@ template < typename std::enable_if_t || std::is_arithmetic_v>* = nullptr> auto allocate_dataframe_buffer( size_t buffer_size, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()) { return rmm::device_uvector(buffer_size, stream_view, mr); @@ -73,7 +73,7 @@ auto allocate_dataframe_buffer( template ::value>* = nullptr> auto allocate_dataframe_buffer( size_t buffer_size, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()) { size_t constexpr tuple_size = cuda::std::tuple_size::value; @@ -83,7 +83,7 @@ auto allocate_dataframe_buffer( template struct dataframe_buffer_type { - using type = decltype(allocate_dataframe_buffer(size_t{0}, rmm::cuda_stream_view{})); + using type = decltype(allocate_dataframe_buffer(size_t{0}, cuda::stream_ref{})); }; template @@ -92,7 +92,7 @@ using dataframe_buffer_type_t = typename dataframe_buffer_type::type; template std::optional> try_allocate_dataframe_buffer( size_t buffer_size, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource_ref()) { try { @@ -134,7 +134,7 @@ using dataframe_buffer_const_iterator_type_t = template void reserve_dataframe_buffer(BufferType& buffer, size_t new_buffer_capacity, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert(is_std_tuple_of_arithmetic_vectors>::value || is_arithmetic_vector, rmm::device_uvector>::value || @@ -151,7 +151,7 @@ void reserve_dataframe_buffer(BufferType& buffer, template void resize_dataframe_buffer(BufferType& buffer, size_t new_buffer_size, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert(is_std_tuple_of_arithmetic_vectors>::value || is_arithmetic_vector, rmm::device_uvector>::value || @@ -166,7 +166,7 @@ void resize_dataframe_buffer(BufferType& buffer, } template -void shrink_to_fit_dataframe_buffer(BufferType& buffer, rmm::cuda_stream_view stream_view) +void shrink_to_fit_dataframe_buffer(BufferType& buffer, cuda::stream_ref stream_view) { static_assert(is_std_tuple_of_arithmetic_vectors>::value || is_arithmetic_vector, rmm::device_uvector>::value || diff --git a/cpp/include/cugraph/utilities/device_comm.hpp b/cpp/include/cugraph/utilities/device_comm.hpp index ad0ddad6a5..02351ef955 100644 --- a/cpp/include/cugraph/utilities/device_comm.hpp +++ b/cpp/include/cugraph/utilities/device_comm.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -11,10 +11,10 @@ #include #include -#include #include #include +#include #include #include @@ -150,7 +150,7 @@ std::enable_if_t::value, void> device_sendre OutputIterator output_first, size_t rx_count, int src, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { // no-op } @@ -166,7 +166,7 @@ device_sendrecv_impl(raft::comms::comms_t const& comm, OutputIterator output_first, size_t rx_count, int src, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { using value_type = typename std::iterator_traits::value_type; static_assert( @@ -177,7 +177,7 @@ device_sendrecv_impl(raft::comms::comms_t const& comm, iter_to_raw_ptr(output_first), rx_count, src, - stream_view.value()); + stream_view.get()); } template @@ -189,7 +189,7 @@ struct device_sendrecv_tuple_iterator_element_impl { OutputIterator output_first, size_t rx_count, int src, - rmm::cuda_stream_view stream_view) const + cuda::stream_ref stream_view) const { using output_value_t = cuda::std::tuple_element_t::value_type>; @@ -203,7 +203,7 @@ struct device_sendrecv_tuple_iterator_element_impl { tuple_element_output_first, rx_count, src, - stream_view.value()); + stream_view.get()); device_sendrecv_tuple_iterator_element_impl().run( comm, input_first, tx_count, dst, output_first, rx_count, src, stream_view); } @@ -218,7 +218,7 @@ struct device_sendrecv_tuple_iterator_element_impl::value, void> device_multic raft::host_span rx_counts, raft::host_span rx_displs, raft::host_span rx_src_ranks, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { // no-op } @@ -252,7 +252,7 @@ device_multicast_sendrecv_impl(raft::comms::comms_t const& comm, raft::host_span rx_counts, raft::host_span rx_displs, raft::host_span rx_src_ranks, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { using value_type = typename std::iterator_traits::value_type; static_assert( @@ -265,7 +265,7 @@ device_multicast_sendrecv_impl(raft::comms::comms_t const& comm, std::vector(rx_counts.begin(), rx_counts.end()), std::vector(rx_displs.begin(), rx_displs.end()), std::vector(rx_src_ranks.begin(), rx_src_ranks.end()), - stream_view.value()); + stream_view.get()); } template @@ -279,7 +279,7 @@ struct device_multicast_sendrecv_tuple_iterator_element_impl { raft::host_span rx_counts, raft::host_span rx_displs, raft::host_span rx_src_ranks, - rmm::cuda_stream_view stream_view) const + cuda::stream_ref stream_view) const { using output_value_t = cuda::std::tuple_element_t::value_type>; @@ -321,7 +321,7 @@ struct device_multicast_sendrecv_tuple_iterator_element_impl rx_counts, raft::host_span rx_displs, raft::host_span rx_src_ranks, - rmm::cuda_stream_view stream_view) const + cuda::stream_ref stream_view) const { } }; @@ -332,7 +332,7 @@ std::enable_if_t::value, void> device_alltoa InputIterator input_first, OutputIterator output_first, size_t count_per_rank, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { // no-op } @@ -345,7 +345,7 @@ device_alltoall_impl(raft::comms::comms_t const& comm, InputIterator input_first, OutputIterator output_first, size_t count_per_rank, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { using value_type = typename std::iterator_traits::value_type; static_assert( @@ -366,12 +366,10 @@ device_alltoall_impl(raft::comms::comms_t const& comm, sizes, displs, ranks, - stream_view.value()); + stream_view.get()); #else - comm.device_alltoall(iter_to_raw_ptr(input_first), - iter_to_raw_ptr(output_first), - count_per_rank, - stream_view.value()); + comm.device_alltoall( + iter_to_raw_ptr(input_first), iter_to_raw_ptr(output_first), count_per_rank, stream_view.get()); #endif } @@ -381,7 +379,7 @@ struct device_alltoall_tuple_iterator_element_impl { InputIterator input_first, OutputIterator output_first, size_t count_per_rank, - rmm::cuda_stream_view stream_view) const + cuda::stream_ref stream_view) const { using output_value_t = typename cuda::std:: tuple_element::value_type>::type; @@ -400,7 +398,7 @@ struct device_alltoall_tuple_iterator_element_impl::value, void> device_bcast_ OutputIterator output_first, size_t count, int root, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { // no-op } @@ -426,12 +424,12 @@ device_bcast_impl(raft::comms::comms_t const& comm, OutputIterator output_first, size_t count, int root, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert(std::is_same_v::value_type, typename std::iterator_traits::value_type>); comm.bcast( - iter_to_raw_ptr(input_first), iter_to_raw_ptr(output_first), count, root, stream_view.value()); + iter_to_raw_ptr(input_first), iter_to_raw_ptr(output_first), count, root, stream_view.get()); } template @@ -441,7 +439,7 @@ struct device_bcast_tuple_iterator_element_impl { OutputIterator output_first, size_t count, int root, - rmm::cuda_stream_view stream_view) const + cuda::stream_ref stream_view) const { device_bcast_impl(comm, cuda::std::get(input_first.get_iterator_tuple()), @@ -461,7 +459,7 @@ struct device_bcast_tuple_iterator_element_impl::value, void> device_allred OutputIterator output_first, size_t count, raft::comms::op_t op, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { // no-op } @@ -487,12 +485,12 @@ device_allreduce_impl(raft::comms::comms_t const& comm, OutputIterator output_first, size_t count, raft::comms::op_t op, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert(std::is_same_v::value_type, typename std::iterator_traits::value_type>); comm.allreduce( - iter_to_raw_ptr(input_first), iter_to_raw_ptr(output_first), count, op, stream_view.value()); + iter_to_raw_ptr(input_first), iter_to_raw_ptr(output_first), count, op, stream_view.get()); } template @@ -502,7 +500,7 @@ struct device_allreduce_tuple_iterator_element_impl { OutputIterator output_first, size_t count, raft::comms::op_t op, - rmm::cuda_stream_view stream_view) const + cuda::stream_ref stream_view) const { device_allreduce_impl(comm, cuda::std::get(input_first.get_iterator_tuple()), @@ -522,7 +520,7 @@ struct device_allreduce_tuple_iterator_element_impl::value, void> device_reduce size_t count, raft::comms::op_t op, int root, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { // no-op } @@ -550,7 +548,7 @@ device_reduce_impl(raft::comms::comms_t const& comm, size_t count, raft::comms::op_t op, int root, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert(std::is_same_v::value_type, typename std::iterator_traits::value_type>); @@ -559,7 +557,7 @@ device_reduce_impl(raft::comms::comms_t const& comm, count, op, root, - stream_view.value()); + stream_view.get()); } template @@ -570,7 +568,7 @@ struct device_reduce_tuple_iterator_element_impl { size_t count, raft::comms::op_t op, int root, - rmm::cuda_stream_view stream_view) const + cuda::stream_ref stream_view) const { device_reduce_impl(comm, cuda::std::get(input_first.get_iterator_tuple()), @@ -592,7 +590,7 @@ struct device_reduce_tuple_iterator_element_impl::value, void> device_allgat InputIterator input_first, OutputIterator output_first, size_t sendcount, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { // no-op } @@ -616,12 +614,12 @@ device_allgather_impl(raft::comms::comms_t const& comm, InputIterator input_first, OutputIterator output_first, size_t sendcount, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert(std::is_same_v::value_type, typename std::iterator_traits::value_type>); comm.allgather( - iter_to_raw_ptr(input_first), iter_to_raw_ptr(output_first), sendcount, stream_view.value()); + iter_to_raw_ptr(input_first), iter_to_raw_ptr(output_first), sendcount, stream_view.get()); } template @@ -630,7 +628,7 @@ struct device_allgather_tuple_iterator_element_impl { InputIterator input_first, OutputIterator output_first, size_t sendcount, - rmm::cuda_stream_view stream_view) const + cuda::stream_ref stream_view) const { device_allgather_impl(comm, cuda::std::get(input_first.get_iterator_tuple()), @@ -648,7 +646,7 @@ struct device_allgather_tuple_iterator_element_impl::value, void> device_allgat OutputIterator output_first, raft::host_span recvcounts, raft::host_span displacements, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { // no-op } @@ -674,7 +672,7 @@ device_allgatherv_impl(raft::comms::comms_t const& comm, OutputIterator output_first, raft::host_span recvcounts, raft::host_span displacements, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert(std::is_same_v::value_type, typename std::iterator_traits::value_type>); @@ -682,7 +680,7 @@ device_allgatherv_impl(raft::comms::comms_t const& comm, iter_to_raw_ptr(output_first), recvcounts.data(), displacements.data(), - stream_view.value()); + stream_view.get()); } template @@ -692,7 +690,7 @@ struct device_allgatherv_tuple_iterator_element_impl { OutputIterator output_first, raft::host_span recvcounts, raft::host_span displacements, - rmm::cuda_stream_view stream_view) const + cuda::stream_ref stream_view) const { device_allgatherv_impl(comm, cuda::std::get(input_first.get_iterator_tuple()), @@ -712,7 +710,7 @@ struct device_allgatherv_tuple_iterator_element_impl recvcounts, raft::host_span displacements, - rmm::cuda_stream_view stream_view) const + cuda::stream_ref stream_view) const { } }; @@ -726,7 +724,7 @@ std::enable_if_t::value, void> device_gather raft::host_span recvcounts, raft::host_span displacements, int root, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { // no-op } @@ -742,7 +740,7 @@ device_gatherv_impl(raft::comms::comms_t const& comm, raft::host_span recvcounts, raft::host_span displacements, int root, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert(std::is_same_v::value_type, typename std::iterator_traits::value_type>); @@ -752,7 +750,7 @@ device_gatherv_impl(raft::comms::comms_t const& comm, recvcounts.data(), displacements.data(), root, - stream_view.value()); + stream_view.get()); } template @@ -764,7 +762,7 @@ struct device_gatherv_tuple_iterator_element_impl { raft::host_span recvcounts, raft::host_span displacements, int root, - rmm::cuda_stream_view stream_view) const + cuda::stream_ref stream_view) const { device_gatherv_impl(comm, cuda::std::get(input_first.get_iterator_tuple()), @@ -788,7 +786,7 @@ struct device_gatherv_tuple_iterator_element_impl recvcounts, raft::host_span displacements, int root, - rmm::cuda_stream_view stream_view) const + cuda::stream_ref stream_view) const { } }; @@ -885,7 +883,7 @@ device_sendrecv(raft::comms::comms_t const& comm, OutputIterator output_first, size_t rx_count, int src, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { detail::device_sendrecv_impl( comm, input_first, tx_count, dst, output_first, rx_count, src, stream_view); @@ -903,7 +901,7 @@ device_sendrecv(raft::comms::comms_t const& comm, OutputIterator output_first, size_t rx_count, int src, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert( cuda::std::tuple_size::value_type>::value == @@ -932,7 +930,7 @@ device_multicast_sendrecv(raft::comms::comms_t const& comm, raft::host_span rx_counts, raft::host_span rx_displs, raft::host_span rx_src_ranks, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { detail::device_multicast_sendrecv_impl(comm, input_first, @@ -960,7 +958,7 @@ device_multicast_sendrecv(raft::comms::comms_t const& comm, raft::host_span rx_counts, raft::host_span rx_displs, raft::host_span rx_src_ranks, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert( cuda::std::tuple_size::value_type>::value == @@ -993,7 +991,7 @@ device_alltoall(raft::comms::comms_t const& comm, InputIterator input_first, OutputIterator output_first, size_t count_per_rank, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { detail::device_alltoall_impl( comm, input_first, output_first, count_per_rank, stream_view); @@ -1008,7 +1006,7 @@ device_alltoall(raft::comms::comms_t const& comm, InputIterator input_first, OutputIterator output_first, size_t count_per_rank, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert( cuda::std::tuple_size::value_type>::value == @@ -1033,7 +1031,7 @@ device_bcast(raft::comms::comms_t const& comm, OutputIterator output_first, size_t count, int root, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { detail::device_bcast_impl(comm, input_first, output_first, count, root, stream_view); } @@ -1048,7 +1046,7 @@ device_bcast(raft::comms::comms_t const& comm, OutputIterator output_first, size_t count, int root, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert( cuda::std::tuple_size::value_type>::value == @@ -1071,7 +1069,7 @@ device_allreduce(raft::comms::comms_t const& comm, OutputIterator output_first, size_t count, raft::comms::op_t op, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { detail::device_allreduce_impl(comm, input_first, output_first, count, op, stream_view); } @@ -1086,7 +1084,7 @@ device_allreduce(raft::comms::comms_t const& comm, OutputIterator output_first, size_t count, raft::comms::op_t op, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert( cuda::std::tuple_size::value_type>::value == @@ -1112,7 +1110,7 @@ device_reduce(raft::comms::comms_t const& comm, size_t count, raft::comms::op_t op, int root, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { detail::device_reduce_impl(comm, input_first, output_first, count, op, root, stream_view); } @@ -1128,7 +1126,7 @@ device_reduce(raft::comms::comms_t const& comm, size_t count, raft::comms::op_t op, int root, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert( cuda::std::tuple_size::value_type>::value == @@ -1152,7 +1150,7 @@ device_allgather(raft::comms::comms_t const& comm, InputIterator input_first, OutputIterator output_first, size_t sendcount, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { detail::device_allgather_impl(comm, input_first, output_first, sendcount, stream_view); } @@ -1166,7 +1164,7 @@ device_allgather(raft::comms::comms_t const& comm, InputIterator input_first, OutputIterator output_first, size_t sendcount, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert( cuda::std::tuple_size::value_type>::value == @@ -1191,7 +1189,7 @@ device_allgatherv(raft::comms::comms_t const& comm, OutputIterator output_first, raft::host_span recvcounts, raft::host_span displacements, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { detail::device_allgatherv_impl( comm, input_first, output_first, recvcounts, displacements, stream_view); @@ -1207,7 +1205,7 @@ device_allgatherv(raft::comms::comms_t const& comm, OutputIterator output_first, raft::host_span recvcounts, raft::host_span displacements, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert( cuda::std::tuple_size::value_type>::value == @@ -1234,7 +1232,7 @@ device_gatherv(raft::comms::comms_t const& comm, raft::host_span recvcounts, raft::host_span displacements, int root, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { detail::device_gatherv_impl( comm, input_first, output_first, sendcount, recvcounts, displacements, root, stream_view); @@ -1252,7 +1250,7 @@ device_gatherv(raft::comms::comms_t const& comm, raft::host_span recvcounts, raft::host_span displacements, int root, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert( cuda::std::tuple_size::value_type>::value == diff --git a/cpp/include/cugraph/utilities/error_check_utils.cuh b/cpp/include/cugraph/utilities/error_check_utils.cuh index 77dcca86d7..6e16e428bc 100644 --- a/cpp/include/cugraph/utilities/error_check_utils.cuh +++ b/cpp/include/cugraph/utilities/error_check_utils.cuh @@ -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 */ diff --git a/cpp/include/cugraph/utilities/groupby_and_count.cuh b/cpp/include/cugraph/utilities/groupby_and_count.cuh index 5ce117ae62..169fadc4cd 100644 --- a/cpp/include/cugraph/utilities/groupby_and_count.cuh +++ b/cpp/include/cugraph/utilities/groupby_and_count.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -22,6 +22,7 @@ #include #include #include +#include #include #include #include @@ -85,7 +86,7 @@ compute_partition_permutation_map(KeyIterator key_first, KeyToGroupIdOp key_to_group_id_op, int group_first, int group_last, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { auto const num_keys = static_cast(cuda::std::distance(key_first, key_last)); auto const num_groups = group_last - group_first; @@ -123,7 +124,7 @@ void apply_multi_partition_permutation( std::tuple, rmm::device_uvector, rmm::device_uvector> const& permutation_map, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { auto const& group_id_offsets = std::get<0>(permutation_map); auto const& intra_partition_displs = std::get<1>(permutation_map); @@ -151,7 +152,7 @@ void multi_partition(ValueIterator value_first, ValueToGroupIdOp value_to_group_id_op, int group_first, int group_last, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { auto permutation_map = compute_partition_permutation_map( value_first, value_last, value_to_group_id_op, group_first, group_last, stream_view); @@ -170,7 +171,7 @@ void multi_partition(KeyIterator key_first, KeyToGroupIdOp key_to_group_id_op, int group_first, int group_last, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { auto const num_keys = static_cast(cuda::std::distance(key_first, key_last)); @@ -187,7 +188,7 @@ template void swap_partitions(ValueIterator value_first, ValueIterator value_last, size_t first_partition_size, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { using value_t = typename thrust::iterator_traits::value_type; @@ -245,7 +246,7 @@ void swap_partitions(KeyIterator key_first, KeyIterator key_last, ValueIterator value_first, size_t first_partition_size, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { using key_t = typename thrust::iterator_traits::value_type; @@ -339,7 +340,7 @@ void mem_frugal_groupby( int num_groups, size_t mem_frugal_threshold, // take the memory frugal approach (instead of thrust::sort) if # // elements to groupby is no smaller than this value - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { CUGRAPH_EXPECTS(!large_buffer_type || large_buffer_manager::memory_buffer_initialized(), @@ -445,7 +446,7 @@ void mem_frugal_groupby( int num_groups, size_t mem_frugal_threshold, // take the memory frugal approach (instead of thrust::sort) if # // elements to groupby is no smaller than this value - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { CUGRAPH_EXPECTS(!large_buffer_type || large_buffer_manager::memory_buffer_initialized(), @@ -563,7 +564,7 @@ rmm::device_uvector groupby_and_count( ValueToGroupIdOp value_to_group_id_op, int num_groups, size_t mem_frugal_threshold, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { CUGRAPH_EXPECTS(!large_buffer_type || large_buffer_manager::memory_buffer_initialized(), @@ -603,7 +604,7 @@ rmm::device_uvector groupby_and_count( KeyToGroupIdOp key_to_group_id_op, int num_groups, size_t mem_frugal_threshold, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { CUGRAPH_EXPECTS(!large_buffer_type || large_buffer_manager::memory_buffer_initialized(), diff --git a/cpp/include/cugraph/utilities/host_scalar_comm.hpp b/cpp/include/cugraph/utilities/host_scalar_comm.hpp index 902a1b6796..bdb919cc39 100644 --- a/cpp/include/cugraph/utilities/host_scalar_comm.hpp +++ b/cpp/include/cugraph/utilities/host_scalar_comm.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -15,6 +15,7 @@ #include #include +#include #include #include @@ -65,12 +66,12 @@ struct host_allreduce_tuple_scalar_element_impl { void run(raft::comms::comms_t const& comm, int64_t* tuple_scalar_elements, raft::comms::op_t op, - cudaStream_t stream) const + cuda::stream_ref stream) const { using element_t = typename cuda::std::tuple_element::type; static_assert(sizeof(element_t) <= sizeof(int64_t)); auto ptr = reinterpret_cast(tuple_scalar_elements + I); - comm.allreduce(ptr, ptr, 1, op, stream); + comm.allreduce(ptr, ptr, 1, op, stream.get()); host_allreduce_tuple_scalar_element_impl().run( comm, tuple_scalar_elements, op, stream); } @@ -81,7 +82,7 @@ struct host_allreduce_tuple_scalar_element_impl { void run(raft::comms::comms_t const& comm, int64_t* tuple_scalar_elements, raft::comms::op_t op, - cudaStream_t stream) const + cuda::stream_ref stream) const { } }; @@ -92,12 +93,12 @@ struct host_reduce_tuple_scalar_element_impl { int64_t* tuple_scalar_elements, raft::comms::op_t op, int root, - cudaStream_t stream) const + cuda::stream_ref stream) const { using element_t = typename cuda::std::tuple_element::type; static_assert(sizeof(element_t) <= sizeof(int64_t)); auto ptr = reinterpret_cast(tuple_scalar_elements + I); - comm.reduce(ptr, ptr, 1, op, root, stream); + comm.reduce(ptr, ptr, 1, op, root, stream.get()); host_reduce_tuple_scalar_element_impl().run( comm, tuple_scalar_elements, op, root, stream); } @@ -109,7 +110,7 @@ struct host_reduce_tuple_scalar_element_impl { int64_t* tuple_scalar_elements, raft::comms::op_t op, int root, - cudaStream_t stream) const + cuda::stream_ref stream) const { } }; @@ -118,7 +119,7 @@ struct host_reduce_tuple_scalar_element_impl { template std::enable_if_t::value, T> host_scalar_allreduce( - raft::comms::comms_t const& comm, T input, raft::comms::op_t op, cudaStream_t stream) + raft::comms::comms_t const& comm, T input, raft::comms::op_t op, cuda::stream_ref stream) { std::variant, rmm::device_uvector> h_tmp_buffer{}; raft::host_span h_tmp_buffer_view{}; @@ -134,9 +135,9 @@ std::enable_if_t::value, T> host_scalar_allreduce( T* d_staging_buffer = d_tmp_buffer.data(); h_staging_buffer[0] = input; raft::update_device(d_staging_buffer, h_staging_buffer, 1, stream); - comm.allreduce(d_staging_buffer, d_staging_buffer, 1, op, stream); + comm.allreduce(d_staging_buffer, d_staging_buffer, 1, op, stream.get()); raft::update_host(h_staging_buffer, d_staging_buffer, 1, stream); - auto status = comm.sync_stream(stream); + auto status = comm.sync_stream(stream.get()); CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); auto h_input = h_staging_buffer[0]; @@ -145,7 +146,7 @@ std::enable_if_t::value, T> host_scalar_allreduce( template std::enable_if_t::value, T> host_scalar_allreduce( - raft::comms::comms_t const& comm, T input, raft::comms::op_t op, cudaStream_t stream) + raft::comms::comms_t const& comm, T input, raft::comms::op_t op, cuda::stream_ref stream) { size_t constexpr tuple_size = cuda::std::tuple_size::value; @@ -169,7 +170,7 @@ std::enable_if_t::value, T> host_scala detail::host_allreduce_tuple_scalar_element_impl().run( comm, d_staging_buffer, op, stream); raft::update_host(h_staging_buffer, d_staging_buffer, tuple_size, stream); - auto status = comm.sync_stream(stream); + auto status = comm.sync_stream(stream.get()); CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); T ret{}; @@ -181,7 +182,11 @@ std::enable_if_t::value, T> host_scala // Return value is valid only in root (return value may better be std::optional in C++17 or later) template std::enable_if_t::value, T> host_scalar_reduce( - raft::comms::comms_t const& comm, T input, raft::comms::op_t op, int root, cudaStream_t stream) + raft::comms::comms_t const& comm, + T input, + raft::comms::op_t op, + int root, + cuda::stream_ref stream) { std::variant, rmm::device_uvector> h_tmp_buffer{}; raft::host_span h_tmp_buffer_view{}; @@ -197,9 +202,9 @@ std::enable_if_t::value, T> host_scalar_reduce( T* d_staging_buffer = d_tmp_buffer.data(); h_staging_buffer[0] = input; raft::update_device(d_staging_buffer, h_staging_buffer, 1, stream); - comm.reduce(d_staging_buffer, d_staging_buffer, 1, op, stream); + comm.reduce(d_staging_buffer, d_staging_buffer, 1, op, stream.get()); if (comm.get_rank() == root) { raft::update_host(h_staging_buffer, d_staging_buffer, 1, stream); } - auto status = comm.sync_stream(stream); + auto status = comm.sync_stream(stream.get()); CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); T h_input{}; @@ -210,7 +215,11 @@ std::enable_if_t::value, T> host_scalar_reduce( // Return value is valid only in root (return value may better be std::optional in C++17 or later) template std::enable_if_t::value, T> host_scalar_reduce( - raft::comms::comms_t const& comm, T input, raft::comms::op_t op, int root, cudaStream_t stream) + raft::comms::comms_t const& comm, + T input, + raft::comms::op_t op, + int root, + cuda::stream_ref stream) { size_t constexpr tuple_size = cuda::std::tuple_size::value; @@ -236,7 +245,7 @@ std::enable_if_t::value, T> host_scala if (comm.get_rank() == root) { raft::update_host(h_staging_buffer, d_staging_buffer, tuple_size, stream); } - auto status = comm.sync_stream(stream); + auto status = comm.sync_stream(stream.get()); CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); T ret{}; @@ -249,7 +258,7 @@ std::enable_if_t::value, T> host_scala template std::enable_if_t::value, T> host_scalar_bcast( - raft::comms::comms_t const& comm, T input, int root, cudaStream_t stream) + raft::comms::comms_t const& comm, T input, int root, cuda::stream_ref stream) { std::variant, rmm::device_uvector> h_tmp_buffer{}; raft::host_span h_tmp_buffer_view{}; @@ -267,9 +276,9 @@ std::enable_if_t::value, T> host_scalar_bcast( h_staging_buffer[0] = input; raft::update_device(d_staging_buffer, h_staging_buffer, 1, stream); } - comm.bcast(d_staging_buffer, 1, root, stream); + comm.bcast(d_staging_buffer, 1, root, stream.get()); if (comm.get_rank() != root) { raft::update_host(h_staging_buffer, d_staging_buffer, 1, stream); } - auto status = comm.sync_stream(stream); + auto status = comm.sync_stream(stream.get()); CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); auto h_input = h_staging_buffer[0]; @@ -278,7 +287,7 @@ std::enable_if_t::value, T> host_scalar_bcast( template std::enable_if_t::value, T> host_scalar_bcast( - raft::comms::comms_t const& comm, T input, int root, cudaStream_t stream) + raft::comms::comms_t const& comm, T input, int root, cuda::stream_ref stream) { size_t constexpr tuple_size = cuda::std::tuple_size::value; @@ -301,11 +310,11 @@ std::enable_if_t::value, T> host_scala .update(h_staging_buffer, input); raft::update_device(d_staging_buffer, h_staging_buffer, tuple_size, stream); } - comm.bcast(d_staging_buffer, tuple_size, root, stream); + comm.bcast(d_staging_buffer, tuple_size, root, stream.get()); if (comm.get_rank() != root) { raft::update_host(h_staging_buffer, d_staging_buffer, tuple_size, stream); } - auto status = comm.sync_stream(stream); + auto status = comm.sync_stream(stream.get()); CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); T ret{}; @@ -316,7 +325,7 @@ std::enable_if_t::value, T> host_scala template std::enable_if_t::value, std::vector> host_scalar_allgather( - raft::comms::comms_t const& comm, T input, cudaStream_t stream) + raft::comms::comms_t const& comm, T input, cuda::stream_ref stream) { std::variant, rmm::device_uvector> h_tmp_buffer{}; raft::host_span h_tmp_buffer_view{}; @@ -333,9 +342,9 @@ std::enable_if_t::value, std::vector> host_scalar_allga h_staging_buffer[comm.get_rank()] = input; raft::update_device( d_staging_buffer + comm.get_rank(), h_staging_buffer + comm.get_rank(), 1, stream); - comm.allgather(d_staging_buffer + comm.get_rank(), d_staging_buffer, size_t{1}, stream); + comm.allgather(d_staging_buffer + comm.get_rank(), d_staging_buffer, size_t{1}, stream.get()); raft::update_host(h_staging_buffer, d_staging_buffer, comm.get_size(), stream); - auto status = comm.sync_stream(stream); + auto status = comm.sync_stream(stream.get()); CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); std::vector h_outputs(h_staging_buffer, h_staging_buffer + comm.get_size()); @@ -344,7 +353,7 @@ std::enable_if_t::value, std::vector> host_scalar_allga template std::enable_if_t::value, std::vector> -host_scalar_allgather(raft::comms::comms_t const& comm, T input, cudaStream_t stream) +host_scalar_allgather(raft::comms::comms_t const& comm, T input, cuda::stream_ref stream) { size_t constexpr tuple_size = cuda::std::tuple_size::value; @@ -369,9 +378,9 @@ host_scalar_allgather(raft::comms::comms_t const& comm, T input, cudaStream_t st tuple_size, stream); comm.allgather( - d_staging_buffer + comm.get_rank() * tuple_size, d_staging_buffer, tuple_size, stream); + d_staging_buffer + comm.get_rank() * tuple_size, d_staging_buffer, tuple_size, stream.get()); raft::update_host(h_staging_buffer, d_staging_buffer, comm.get_size() * tuple_size, stream); - auto status = comm.sync_stream(stream); + auto status = comm.sync_stream(stream.get()); CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); std::vector ret(comm.get_size()); @@ -387,7 +396,7 @@ std::enable_if_t::value, T> host_scalar_scatter( raft::comms::comms_t const& comm, std::vector const& inputs, // relevant only in root int root, - cudaStream_t stream) + cuda::stream_ref stream) { CUGRAPH_EXPECTS( ((comm.get_rank() == root) && (inputs.size() == static_cast(comm.get_size()))) || @@ -410,12 +419,12 @@ std::enable_if_t::value, T> host_scalar_scatter( std::copy(inputs.begin(), inputs.end(), h_staging_buffer); raft::update_device(d_staging_buffer, h_staging_buffer, comm.get_size(), stream); } - comm.bcast(d_staging_buffer, comm.get_size(), root, stream); + comm.bcast(d_staging_buffer, comm.get_size(), root, stream.get()); if (comm.get_rank() != root) { raft::update_host( h_staging_buffer + comm.get_rank(), d_staging_buffer + comm.get_rank(), 1, stream); } - auto status = comm.sync_stream(stream); + auto status = comm.sync_stream(stream.get()); CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); auto h_output = h_staging_buffer[comm.get_rank()]; @@ -427,7 +436,7 @@ std::enable_if_t::value, T> host_scala raft::comms::comms_t const& comm, std::vector const& inputs, // relevant only in root int root, - cudaStream_t stream) + cuda::stream_ref stream) { size_t constexpr tuple_size = cuda::std::tuple_size::value; CUGRAPH_EXPECTS( @@ -456,14 +465,14 @@ std::enable_if_t::value, T> host_scala } raft::update_device(d_staging_buffer, h_staging_buffer, comm.get_size() * tuple_size, stream); } - comm.bcast(d_staging_buffer, comm.get_size() * tuple_size, root, stream); + comm.bcast(d_staging_buffer, comm.get_size() * tuple_size, root, stream.get()); if (comm.get_rank() != root) { raft::update_host(h_staging_buffer + comm.get_rank() * tuple_size, d_staging_buffer + comm.get_rank() * tuple_size, tuple_size, stream); } - auto status = comm.sync_stream(stream); + auto status = comm.sync_stream(stream.get()); CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); T ret{}; @@ -476,7 +485,7 @@ std::enable_if_t::value, T> host_scala // Return value is valid only in root (return value may better be std::optional in C++17 or later) template std::enable_if_t::value, std::vector> host_scalar_gather( - raft::comms::comms_t const& comm, T input, int root, cudaStream_t stream) + raft::comms::comms_t const& comm, T input, int root, cuda::stream_ref stream) { std::variant, rmm::device_uvector> h_tmp_buffer{}; raft::host_span h_tmp_buffer_view{}; @@ -493,11 +502,11 @@ std::enable_if_t::value, std::vector> host_scalar_gathe h_staging_buffer[comm.get_rank()] = input; raft::update_device( d_staging_buffer + comm.get_rank(), h_staging_buffer + comm.get_rank(), 1, stream); - comm.gather(d_staging_buffer + comm.get_rank(), d_staging_buffer, size_t{1}, root, stream); + comm.gather(d_staging_buffer + comm.get_rank(), d_staging_buffer, size_t{1}, root, stream.get()); if (comm.get_rank() == root) { raft::update_host(h_staging_buffer, d_staging_buffer, comm.get_size(), stream); } - auto status = comm.sync_stream(stream); + auto status = comm.sync_stream(stream.get()); CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); std::vector h_outputs{}; @@ -510,7 +519,7 @@ std::enable_if_t::value, std::vector> host_scalar_gathe // Return value is valid only in root (return value may better be std::optional in C++17 or later) template std::enable_if_t::value, std::vector> -host_scalar_gather(raft::comms::comms_t const& comm, T input, int root, cudaStream_t stream) +host_scalar_gather(raft::comms::comms_t const& comm, T input, int root, cuda::stream_ref stream) { size_t constexpr tuple_size = cuda::std::tuple_size::value; @@ -534,12 +543,15 @@ host_scalar_gather(raft::comms::comms_t const& comm, T input, int root, cudaStre h_staging_buffer + comm.get_rank() * tuple_size, tuple_size, stream); - comm.gather( - d_staging_buffer + comm.get_rank() * tuple_size, d_staging_buffer, tuple_size, root, stream); + comm.gather(d_staging_buffer + comm.get_rank() * tuple_size, + d_staging_buffer, + tuple_size, + root, + stream.get()); if (comm.get_rank() == root) { raft::update_host(h_staging_buffer, d_staging_buffer, comm.get_size() * tuple_size, stream); } - auto status = comm.sync_stream(stream); + auto status = comm.sync_stream(stream.get()); CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); std::vector ret(comm.get_size()); diff --git a/cpp/include/cugraph/utilities/mask_utils.cuh b/cpp/include/cugraph/utilities/mask_utils.cuh index df873b13e8..f497531cfa 100644 --- a/cpp/include/cugraph/utilities/mask_utils.cuh +++ b/cpp/include/cugraph/utilities/mask_utils.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 */ #pragma once @@ -17,6 +17,7 @@ #include #include #include +#include #include #include #include @@ -341,7 +342,7 @@ void partition_by_mask(InputIterator input_first, uint32_t const* mask_first, size_t first_size, size_t second_size, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { using element_t = typename thrust::iterator_traits::value_type; @@ -372,7 +373,7 @@ struct partition_by_mask_zip_split_impl { uint32_t const* mask_first, size_t first_size, size_t second_size, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type) { auto const& input_tuple = input_first.get_iterator_tuple(); @@ -397,7 +398,7 @@ struct partition_by_mask_zip_split_impl { uint32_t const*, size_t, size_t, - rmm::cuda_stream_view, + cuda::stream_ref, std::optional) { } @@ -410,7 +411,7 @@ void partition_by_mask_zip_split( uint32_t const* mask_first, size_t first_size, size_t second_size, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { constexpr size_t tuple_size = cuda::std::tuple_size::value; @@ -427,7 +428,7 @@ void partition_by_mask(ZipIterator input_first, uint32_t const* mask_first, size_t first_size, size_t second_size, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { partition_by_mask_zip_split( @@ -499,7 +500,7 @@ template std::tuple> mark_entries( size_t num_entries, comparison_t comparison, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { auto marked_entries = large_buffer_type ? large_buffer_manager::allocate_memory_buffer( @@ -556,7 +557,7 @@ void partition_by_mask(InputIterator input_first, uint32_t const* mask_first, size_t first_size, size_t second_size, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type) { detail::partition_by_mask( diff --git a/cpp/include/cugraph/utilities/mem_frugal_partition.cuh b/cpp/include/cugraph/utilities/mem_frugal_partition.cuh index 16e2af7f5e..b657fcecbd 100644 --- a/cpp/include/cugraph/utilities/mem_frugal_partition.cuh +++ b/cpp/include/cugraph/utilities/mem_frugal_partition.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -12,6 +12,7 @@ #include #include +#include namespace CUGRAPH_EXPORT cugraph { @@ -27,7 +28,7 @@ ValueIterator mem_frugal_partition( ValueIterator value_last, ValueToGroupIdOp value_to_group_id_op, int pivot, // group id less than pivot goes to the first partition - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { CUGRAPH_EXPECTS(!large_buffer_type || large_buffer_manager::memory_buffer_initialized(), @@ -64,7 +65,7 @@ std::tuple mem_frugal_partition( ValueIterator value_first, KeyToGroupIdOp key_to_group_id_op, int pivot, // group Id less than pivot goes to the first partition - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { CUGRAPH_EXPECTS(!large_buffer_type || large_buffer_manager::memory_buffer_initialized(), diff --git a/cpp/include/cugraph/utilities/misc_utils.cuh b/cpp/include/cugraph/utilities/misc_utils.cuh index a4d843c024..06b76f5cb7 100644 --- a/cpp/include/cugraph/utilities/misc_utils.cuh +++ b/cpp/include/cugraph/utilities/misc_utils.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -20,6 +20,7 @@ #include #include #include +#include #include #include #include @@ -101,7 +102,7 @@ std::optional to_std_optional(cuda::std::optional val) template rmm::device_uvector expand_sparse_offsets(raft::device_span offsets, idx_t base_idx, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { assert(offsets.size() > 0); diff --git a/cpp/include/cugraph/utilities/partition_scatter_map_wrappers.cuh b/cpp/include/cugraph/utilities/partition_scatter_map_wrappers.cuh index e401c701e6..8c5db5d26f 100644 --- a/cpp/include/cugraph/utilities/partition_scatter_map_wrappers.cuh +++ b/cpp/include/cugraph/utilities/partition_scatter_map_wrappers.cuh @@ -1,14 +1,15 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once #include -#include #include +#include + #include #include #include @@ -27,14 +28,14 @@ CUGRAPH_EXPORT rmm::device_uvector compute_partition_scatter_map_impl( offset_t const* intra_partition_displs, size_t const* group_displacements, size_t num_elements, - rmm::cuda_stream_view stream_view); + cuda::stream_ref stream_view); template rmm::device_uvector compute_partition_scatter_map( rmm::device_uvector const& group_id_offsets, rmm::device_uvector const& intra_partition_displs, rmm::device_uvector const& group_displacements, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { static_assert(compute_partition_scatter_map_supported_v, "compute_partition_scatter_map is not explicitly instantiated for this " diff --git a/cpp/include/cugraph/utilities/permute_wrappers.cuh b/cpp/include/cugraph/utilities/permute_wrappers.cuh index d160dc9e34..d402eb1d7a 100644 --- a/cpp/include/cugraph/utilities/permute_wrappers.cuh +++ b/cpp/include/cugraph/utilities/permute_wrappers.cuh @@ -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 */ #pragma once @@ -9,10 +9,10 @@ #include #include -#include #include #include +#include #include #include @@ -32,13 +32,13 @@ template void permute_in_place_impl(T* first, std::size_t const* map_first, std::size_t num_elements, - rmm::cuda_stream_view stream_view); + cuda::stream_ref stream_view); template void permute_scalar_in_place(Iterator first, std::size_t const* map_first, std::size_t num_elements, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { if constexpr (scatter_supported_scalar_value_v>) { permute_in_place_impl(first, map_first, num_elements, stream_view); @@ -58,7 +58,7 @@ struct permute_zip_in_place_split_impl { static void run(ZipIterator first, ZipIterator last, std::size_t const* map_first, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { auto const num_elements = static_cast(cuda::std::distance(first, last)); permute_scalar_in_place( @@ -70,12 +70,14 @@ struct permute_zip_in_place_split_impl { template struct permute_zip_in_place_split_impl { - static void run(ZipIterator, ZipIterator, std::size_t const*, rmm::cuda_stream_view) {} + static void run(ZipIterator, ZipIterator, std::size_t const*, cuda::stream_ref) {} }; template -std::enable_if_t> permute_in_place( - Iterator first, Iterator last, std::size_t const* map_first, rmm::cuda_stream_view stream_view) +std::enable_if_t> permute_in_place(Iterator first, + Iterator last, + std::size_t const* map_first, + cuda::stream_ref stream_view) { permute_zip_in_place_split_impl>::run( first, last, map_first, stream_view); @@ -83,7 +85,7 @@ std::enable_if_t> permute_in_place( template std::enable_if_t> permute_in_place( - Iterator first, Iterator last, std::size_t const* map_first, rmm::cuda_stream_view stream_view) + Iterator first, Iterator last, std::size_t const* map_first, cuda::stream_ref stream_view) { auto const num_elements = static_cast(cuda::std::distance(first, last)); permute_scalar_in_place(first, map_first, num_elements, stream_view); diff --git a/cpp/include/cugraph/utilities/shuffle_comm.cuh b/cpp/include/cugraph/utilities/shuffle_comm.cuh index a2078c1a37..7b570b852c 100644 --- a/cpp/include/cugraph/utilities/shuffle_comm.cuh +++ b/cpp/include/cugraph/utilities/shuffle_comm.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -23,6 +23,7 @@ #include #include #include +#include #include #include #include @@ -56,7 +57,7 @@ inline std::tuple, compute_tx_rx_counts_displs_ranks(raft::comms::comms_t const& comm, raft::device_span d_tx_value_counts, bool drop_empty_ranks, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { auto const comm_size = comm.get_size(); @@ -74,9 +75,9 @@ compute_tx_rx_counts_displs_ranks(raft::comms::comms_t const& comm, rmm::device_uvector d_rx_value_counts(comm_size, stream_view); device_alltoall(comm, d_tx_value_counts.data(), d_rx_value_counts.data(), size_t{1}, stream_view); - raft::update_host(tx_counts.data(), d_tx_value_counts.data(), comm_size, stream_view.value()); - raft::update_host(rx_counts.data(), d_rx_value_counts.data(), comm_size, stream_view.value()); - stream_view.synchronize(); + raft::update_host(tx_counts.data(), d_tx_value_counts.data(), comm_size, stream_view); + raft::update_host(rx_counts.data(), d_rx_value_counts.data(), comm_size, stream_view); + stream_view.sync(); std::partial_sum(tx_counts.begin(), tx_counts.end() - 1, tx_displs.begin() + 1); std::partial_sum(rx_counts.begin(), rx_counts.end() - 1, rx_displs.begin() + 1); @@ -139,7 +140,7 @@ template auto shuffle_values(raft::comms::comms_t const& comm, TxValueIterator tx_value_first, raft::device_span d_tx_value_counts, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { using value_t = typename thrust::iterator_traits::value_type; @@ -188,7 +189,7 @@ template auto shuffle_values(raft::comms::comms_t const& comm, TxValueIterator tx_value_first, raft::host_span tx_value_counts, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { using value_t = typename thrust::iterator_traits::value_type; @@ -197,7 +198,7 @@ auto shuffle_values(raft::comms::comms_t const& comm, rmm::device_uvector d_tx_value_counts(comm_size, stream_view); raft::update_device( - d_tx_value_counts.data(), tx_value_counts.data(), comm_size, stream_view.value()); + d_tx_value_counts.data(), tx_value_counts.data(), comm_size, stream_view.get()); return shuffle_values( comm, @@ -217,7 +218,7 @@ auto shuffle_values( raft::host_span tx_value_counts, size_t alignment, // # elements std::optional::value_type> fill_value, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { using value_t = typename thrust::iterator_traits::value_type; @@ -296,7 +297,7 @@ auto shuffle_values( stream_view); raft::update_host( rx_aligned_counts.data(), d_rx_aligned_counts.data(), d_rx_aligned_counts.size(), stream_view); - RAFT_CUDA_TRY(cudaStreamSynchronize(stream_view)); + stream_view.sync(); size_t offset{0}; for (size_t i = 0; i < rx_counts.size(); ++i) { auto target_alignment = (alignment - rx_unaligned_counts[i]) % alignment; @@ -364,7 +365,7 @@ auto shuffle_and_unique_segment_sorted_values( // tx_value_counts[i], where i = [0, comm_size); and bettter be // unique to reduce communication volume raft::host_span tx_value_counts, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { using value_t = typename thrust::iterator_traits::value_type; @@ -394,7 +395,7 @@ auto shuffle_and_unique_segment_sorted_values( } else { rmm::device_uvector d_tx_value_counts(comm_size, stream_view); raft::update_device( - d_tx_value_counts.data(), tx_value_counts.data(), comm_size, stream_view.value()); + d_tx_value_counts.data(), tx_value_counts.data(), comm_size, stream_view.get()); std::vector tx_counts{}; std::vector tx_displs{}; @@ -469,7 +470,7 @@ auto groupby_gpu_id_and_shuffle_values( ValueIterator tx_value_first /* [INOUT */, ValueIterator tx_value_last /* [INOUT */, ValueToGPUIdOp value_to_gpu_id_op, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { using value_t = typename thrust::iterator_traits::value_type; @@ -536,7 +537,7 @@ auto groupby_gpu_id_and_shuffle_kv_pairs( KeyIterator tx_key_last /* [INOUT */, ValueIterator tx_value_first /* [INOUT */, KeyToGPUIdOp key_to_gpu_id_op, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, std::optional large_buffer_type = std::nullopt) { using key_t = typename thrust::iterator_traits::value_type; diff --git a/cpp/libcugraph_etl/include/hash/concurrent_unordered_map.cuh b/cpp/libcugraph_etl/include/hash/concurrent_unordered_map.cuh index efe17d3544..d2e8097373 100644 --- a/cpp/libcugraph_etl/include/hash/concurrent_unordered_map.cuh +++ b/cpp/libcugraph_etl/include/hash/concurrent_unordered_map.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2017-2026, NVIDIA CORPORATION. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2017-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -13,12 +13,12 @@ #include #include -#include #include #include #include #include +#include #include #include @@ -153,7 +153,7 @@ class concurrent_unordered_map { * storage */ static auto create(size_type capacity, - rmm::cuda_stream_view stream, + cuda::stream_ref stream, mapped_type const unused_element = std::numeric_limits::max(), key_type const unused_key = std::numeric_limits::max(), Hasher const& hash_function = hasher(), @@ -423,7 +423,7 @@ class concurrent_unordered_map { } } - void assign_async(concurrent_unordered_map const& other, rmm::cuda_stream_view stream) + void assign_async(concurrent_unordered_map const& other, cuda::stream_ref stream) { if (other.m_capacity <= m_capacity) { m_capacity = other.m_capacity; @@ -438,13 +438,13 @@ class concurrent_unordered_map { other.m_hashtbl_values, m_capacity * sizeof(value_type), cudaMemcpyDefault, - stream.value())); + stream.get())); } - void clear_async(rmm::cuda_stream_view stream) + void clear_async(cuda::stream_ref stream) { constexpr int block_size = 128; - init_hashtbl<<<((m_capacity - 1) / block_size) + 1, block_size, 0, stream.value()>>>( + init_hashtbl<<<((m_capacity - 1) / block_size) + 1, block_size, 0, stream.get()>>>( m_hashtbl_values, m_capacity, m_unused_key, m_unused_element); } @@ -456,7 +456,7 @@ class concurrent_unordered_map { } } - void prefetch(int const dev_id, rmm::cuda_stream_view stream) + void prefetch(int const dev_id, cuda::stream_ref stream) { cudaPointerAttributes hashtbl_values_ptr_attributes; cudaError_t status = cudaPointerGetAttributes(&hashtbl_values_ptr_attributes, m_hashtbl_values); @@ -476,7 +476,7 @@ class concurrent_unordered_map { * * @param stream CUDA stream used for device memory operations and kernel launches. */ - void destroy(rmm::cuda_stream_view stream) + void destroy(cuda::stream_ref stream) { m_allocator.deallocate(m_hashtbl_values, m_capacity, stream); delete this; @@ -517,7 +517,7 @@ class concurrent_unordered_map { Hasher const& hash_function, Equality const& equal, allocator_type const& allocator, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) : m_hf(hash_function), m_equal(equal), m_allocator(allocator), @@ -541,10 +541,10 @@ class concurrent_unordered_map { } if (m_capacity > 0) { - init_hashtbl<<<((m_capacity - 1) / block_size) + 1, block_size, 0, stream.value()>>>( + init_hashtbl<<<((m_capacity - 1) / block_size) + 1, block_size, 0, stream.get()>>>( m_hashtbl_values, m_capacity, m_unused_key, m_unused_element); } - CUDF_CHECK_CUDA(stream.value()); + CUDF_CHECK_CUDA(stream.get()); } }; diff --git a/cpp/libcugraph_etl/src/renumbering.cu b/cpp/libcugraph_etl/src/renumbering.cu index b628c55b67..769c5bfc38 100644 --- a/cpp/libcugraph_etl/src/renumbering.cu +++ b/cpp/libcugraph_etl/src/renumbering.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include @@ -778,7 +778,7 @@ struct renumber_functor { const_cast(str_col_view.offsets().data())); } - cudaStream_t exec_strm = handle.get_stream(); + cudaStream_t exec_strm = handle.get_stream().get(); auto mr = rmm::mr::pinned_host_memory_resource(); size_t hist_size = sizeof(accum_type) * 32; diff --git a/cpp/src/c_api/array.hpp b/cpp/src/c_api/array.hpp index 95e54bd135..a2e1f26920 100644 --- a/cpp/src/c_api/array.hpp +++ b/cpp/src/c_api/array.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -9,6 +9,8 @@ #include #include +#include + #include namespace cugraph { @@ -44,7 +46,7 @@ struct cugraph_type_erased_device_array_t { cugraph_type_erased_device_array_t(size_t size, size_t num_bytes, cugraph_data_type_id_t type, - rmm::cuda_stream_view const& stream_view) + cuda::stream_ref const& stream_view) : size_(size), data_(num_bytes, stream_view), type_(type) { } diff --git a/cpp/src/c_api/bfs.cpp b/cpp/src/c_api/bfs.cpp index 26a4eeb83b..63b121d644 100644 --- a/cpp/src/c_api/bfs.cpp +++ b/cpp/src/c_api/bfs.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/src/c_api/extract_ego.cpp b/cpp/src/c_api/extract_ego.cpp index c79eefb771..f4b9dc4e72 100644 --- a/cpp/src/c_api/extract_ego.cpp +++ b/cpp/src/c_api/extract_ego.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/src/c_api/graph_generators.cpp b/cpp/src/c_api/graph_generators.cpp index c57ad45424..66d3f3b6e2 100644 --- a/cpp/src/c_api/graph_generators.cpp +++ b/cpp/src/c_api/graph_generators.cpp @@ -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 */ #include "c_api/array.hpp" diff --git a/cpp/src/c_api/graph_helper.hpp b/cpp/src/c_api/graph_helper.hpp index 936984fab9..4e9ee29ed1 100644 --- a/cpp/src/c_api/graph_helper.hpp +++ b/cpp/src/c_api/graph_helper.hpp @@ -7,13 +7,15 @@ #include +#include + namespace cugraph { namespace c_api { template rmm::device_uvector expand_sparse_offsets(raft::device_span offsets, vertex_t base_vertex_id, - rmm::cuda_stream_view const& stream); + cuda::stream_ref const& stream); template edge_property_t create_constant_edge_property( @@ -34,7 +36,7 @@ edge_property_t create_constant_edge_prope template void copy_or_transform(raft::device_span output, cugraph_type_erased_device_array_view_t const* input, - rmm::cuda_stream_view const& stream_view); + cuda::stream_ref const& stream_view); } // namespace c_api } // namespace cugraph diff --git a/cpp/src/c_api/graph_helper_impl.cuh b/cpp/src/c_api/graph_helper_impl.cuh index b10b8601d8..d4923e58ad 100644 --- a/cpp/src/c_api/graph_helper_impl.cuh +++ b/cpp/src/c_api/graph_helper_impl.cuh @@ -10,6 +10,7 @@ #include #include +#include namespace cugraph { namespace c_api { @@ -17,7 +18,7 @@ namespace c_api { template rmm::device_uvector expand_sparse_offsets(raft::device_span offsets, vertex_t base_vertex_id, - rmm::cuda_stream_view const& stream) + cuda::stream_ref const& stream) { return cugraph::detail::expand_sparse_offsets(offsets, base_vertex_id, stream); } @@ -36,7 +37,7 @@ edge_property_t create_constant_edge_prope template void copy_or_transform(raft::device_span output, cugraph_type_erased_device_array_view_t const* input, - rmm::cuda_stream_view const& stream_view) + cuda::stream_ref const& stream_view) { if (((input->type_ == cugraph_data_type_id_t::INT8) && (std::is_same_v)) || ((input->type_ == cugraph_data_type_id_t::INT16) && (std::is_same_v)) || diff --git a/cpp/src/c_api/graph_helper_sg.cu b/cpp/src/c_api/graph_helper_sg.cu index be93fdc9f1..948d6ce876 100644 --- a/cpp/src/c_api/graph_helper_sg.cu +++ b/cpp/src/c_api/graph_helper_sg.cu @@ -7,61 +7,53 @@ #include +#include + namespace cugraph { namespace c_api { template CUGRAPH_EXPORT rmm::device_uvector expand_sparse_offsets( - raft::device_span offsets, - int32_t base_vertex_id, - rmm::cuda_stream_view const& stream); + raft::device_span offsets, int32_t base_vertex_id, cuda::stream_ref const& stream); template CUGRAPH_EXPORT rmm::device_uvector expand_sparse_offsets( - raft::device_span offsets, - int32_t base_vertex_id, - rmm::cuda_stream_view const& stream); + raft::device_span offsets, int32_t base_vertex_id, cuda::stream_ref const& stream); template CUGRAPH_EXPORT rmm::device_uvector expand_sparse_offsets( - raft::device_span offsets, - int64_t base_vertex_id, - rmm::cuda_stream_view const& stream); + raft::device_span offsets, int64_t base_vertex_id, cuda::stream_ref const& stream); template CUGRAPH_EXPORT rmm::device_uvector expand_sparse_offsets( - raft::device_span offsets, - int32_t base_vertex_id, - rmm::cuda_stream_view const& stream); + raft::device_span offsets, int32_t base_vertex_id, cuda::stream_ref const& stream); template CUGRAPH_EXPORT rmm::device_uvector expand_sparse_offsets( - raft::device_span offsets, - int64_t base_vertex_id, - rmm::cuda_stream_view const& stream); + raft::device_span offsets, int64_t base_vertex_id, cuda::stream_ref const& stream); template CUGRAPH_EXPORT void copy_or_transform(raft::device_span output, cugraph_type_erased_device_array_view_t const* input, - rmm::cuda_stream_view const& stream_view); + cuda::stream_ref const& stream_view); template CUGRAPH_EXPORT void copy_or_transform(raft::device_span output, cugraph_type_erased_device_array_view_t const* input, - rmm::cuda_stream_view const& stream_view); + cuda::stream_ref const& stream_view); template CUGRAPH_EXPORT void copy_or_transform(raft::device_span output, cugraph_type_erased_device_array_view_t const* input, - rmm::cuda_stream_view const& stream_view); + cuda::stream_ref const& stream_view); template CUGRAPH_EXPORT void copy_or_transform(raft::device_span output, cugraph_type_erased_device_array_view_t const* input, - rmm::cuda_stream_view const& stream_view); + cuda::stream_ref const& stream_view); template CUGRAPH_EXPORT void copy_or_transform(raft::device_span output, cugraph_type_erased_device_array_view_t const* input, - rmm::cuda_stream_view const& stream_view); + cuda::stream_ref const& stream_view); template CUGRAPH_EXPORT void copy_or_transform(raft::device_span output, cugraph_type_erased_device_array_view_t const* input, - rmm::cuda_stream_view const& stream_view); + cuda::stream_ref const& stream_view); template CUGRAPH_EXPORT void copy_or_transform(raft::device_span output, cugraph_type_erased_device_array_view_t const* input, - rmm::cuda_stream_view const& stream_view); + cuda::stream_ref const& stream_view); template CUGRAPH_EXPORT edge_property_t create_constant_edge_property( raft::handle_t const& handle, diff --git a/cpp/src/c_api/graph_mg.cpp b/cpp/src/c_api/graph_mg.cpp index 98b6ce6459..b5674123df 100644 --- a/cpp/src/c_api/graph_mg.cpp +++ b/cpp/src/c_api/graph_mg.cpp @@ -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 */ diff --git a/cpp/src/c_api/random.cpp b/cpp/src/c_api/random.cpp index 2eeb5f2ffd..fbaf4e9a70 100644 --- a/cpp/src/c_api/random.cpp +++ b/cpp/src/c_api/random.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/src/centrality/betweenness_centrality_impl.cuh b/cpp/src/centrality/betweenness_centrality_impl.cuh index 7f31a8d9e1..74718148e9 100644 --- a/cpp/src/centrality/betweenness_centrality_impl.cuh +++ b/cpp/src/centrality/betweenness_centrality_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -1043,7 +1043,7 @@ void multisource_backward_pass( num_segments_in_chunk, offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); if (temp_storage_bytes > d_tmp_storage.size()) { d_tmp_storage = rmm::device_uvector(temp_storage_bytes, handle.get_stream()); @@ -1059,7 +1059,7 @@ void multisource_backward_pass( num_segments_in_chunk, offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); } } diff --git a/cpp/src/community/approx_weighted_matching_impl.cuh b/cpp/src/community/approx_weighted_matching_impl.cuh index 86cda764e2..4ab1690902 100644 --- a/cpp/src/community/approx_weighted_matching_impl.cuh +++ b/cpp/src/community/approx_weighted_matching_impl.cuh @@ -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 */ #pragma once diff --git a/cpp/src/community/detail/common_methods.cuh b/cpp/src/community/detail/common_methods.cuh index e325c384e7..6b539d3cb9 100644 --- a/cpp/src/community/detail/common_methods.cuh +++ b/cpp/src/community/detail/common_methods.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/src/community/detail/common_methods.hpp b/cpp/src/community/detail/common_methods.hpp index 34351c9f57..5d60850035 100644 --- a/cpp/src/community/detail/common_methods.hpp +++ b/cpp/src/community/detail/common_methods.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -37,11 +37,11 @@ void timer_stop(raft::handle_t const& handle, HighResTimer& hr_timer) { if constexpr (multi_gpu) { if (handle.get_comms().get_rank() == 0) { - handle.get_stream().synchronize(); + handle.get_stream().sync(); hr_timer.stop(); } } else { - handle.get_stream().synchronize(); + handle.get_stream().sync(); hr_timer.stop(); } } diff --git a/cpp/src/community/detail/maximal_independent_moves.cuh b/cpp/src/community/detail/maximal_independent_moves.cuh index 710e0dc100..4daba99fa0 100644 --- a/cpp/src/community/detail/maximal_independent_moves.cuh +++ b/cpp/src/community/detail/maximal_independent_moves.cuh @@ -1,6 +1,6 @@ /* - * 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 diff --git a/cpp/src/community/detail/refine_impl.cuh b/cpp/src/community/detail/refine_impl.cuh index 8e0938f1f9..d964505495 100644 --- a/cpp/src/community/detail/refine_impl.cuh +++ b/cpp/src/community/detail/refine_impl.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 */ #pragma once diff --git a/cpp/src/community/edge_triangle_count_impl.cuh b/cpp/src/community/edge_triangle_count_impl.cuh index 672b22fbfc..d996376b21 100644 --- a/cpp/src/community/edge_triangle_count_impl.cuh +++ b/cpp/src/community/edge_triangle_count_impl.cuh @@ -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 */ diff --git a/cpp/src/community/egonet_impl.cuh b/cpp/src/community/egonet_impl.cuh index 6dbdafe093..514a6cf4e4 100644 --- a/cpp/src/community/egonet_impl.cuh +++ b/cpp/src/community/egonet_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -17,7 +17,6 @@ #endif #include -#include #include #include @@ -83,7 +82,7 @@ extract(raft::handle_t const& handle, std::vector> reached{}; reached.reserve(num_sources); - user_stream_view.synchronize(); + user_stream_view.sync(); #ifdef TIMING HighResTimer hr_timer; hr_timer.start("ego_neighbors"); @@ -165,9 +164,9 @@ extract(raft::handle_t const& handle, h_neighbors_offsets[i + 1] = h_neighbors_offsets[i] + reached[i].size(); } raft::update_device( - neighbors_offsets.data(), &h_neighbors_offsets[0], num_sources + 1, user_stream_view.value()); - neighbors.resize(h_neighbors_offsets[num_sources], user_stream_view.value()); - user_stream_view.synchronize(); + neighbors_offsets.data(), &h_neighbors_offsets[0], num_sources + 1, user_stream_view.get()); + neighbors.resize(h_neighbors_offsets[num_sources], user_stream_view.get()); + user_stream_view.sync(); // Construct the neighbors list concurrently for (size_t i = 0; i < num_sources; i++) { diff --git a/cpp/src/community/legacy/spectral_clustering.cu b/cpp/src/community/legacy/spectral_clustering.cu index 9789044a3f..538eeeb756 100644 --- a/cpp/src/community/legacy/spectral_clustering.cu +++ b/cpp/src/community/legacy/spectral_clustering.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -76,7 +76,7 @@ void balancedCutClustering_impl(raft::handle_t const& handle, static_cast(graph.number_of_vertices), src_indices.data(), static_cast(graph.number_of_edges), - handle.get_stream()); + handle.get_stream().get()); // Create coordinate structure view from converted COO data auto coord_view = raft::make_device_coordinate_structure_view( @@ -151,7 +151,7 @@ void spectralModularityMaximization_impl( static_cast(graph.number_of_vertices), src_indices.data(), static_cast(graph.number_of_edges), - handle.get_stream()); + handle.get_stream().get()); // Create coordinate structure view from converted COO data auto coord_view = raft::make_device_coordinate_structure_view( diff --git a/cpp/src/community/leiden_impl.cuh b/cpp/src/community/leiden_impl.cuh index 78c274f9e3..c30d15081d 100644 --- a/cpp/src/community/leiden_impl.cuh +++ b/cpp/src/community/leiden_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/src/community/triangle_count_impl.cuh b/cpp/src/community/triangle_count_impl.cuh index a08977efb4..6883bbd99c 100644 --- a/cpp/src/community/triangle_count_impl.cuh +++ b/cpp/src/community/triangle_count_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/src/components/mis_impl.cuh b/cpp/src/components/mis_impl.cuh index 7bf7d26a77..a48ca07703 100644 --- a/cpp/src/components/mis_impl.cuh +++ b/cpp/src/components/mis_impl.cuh @@ -1,6 +1,6 @@ /* - * 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 diff --git a/cpp/src/components/strongly_connected_components_impl.cuh b/cpp/src/components/strongly_connected_components_impl.cuh index bb26b8cde8..6a8a540db6 100644 --- a/cpp/src/components/strongly_connected_components_impl.cuh +++ b/cpp/src/components/strongly_connected_components_impl.cuh @@ -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 */ #pragma once diff --git a/cpp/src/converters/legacy/COOtoCSR.cuh b/cpp/src/converters/legacy/COOtoCSR.cuh index 0f8f0bd653..b82c440e64 100644 --- a/cpp/src/converters/legacy/COOtoCSR.cuh +++ b/cpp/src/converters/legacy/COOtoCSR.cuh @@ -25,6 +25,7 @@ #include #include #include +#include #include #include #include @@ -56,7 +57,7 @@ namespace detail { * @param[out] result Total number of vertices */ template -VT sort(legacy::GraphCOOView& graph, rmm::cuda_stream_view stream_view) +VT sort(legacy::GraphCOOView& graph, cuda::stream_ref stream_view) { VT max_src_id; VT max_dst_id; @@ -92,11 +93,8 @@ VT sort(legacy::GraphCOOView& graph, rmm::cuda_stream_view stream_vi } template -void fill_offset(VT* source, - ET* offsets, - VT number_of_vertices, - ET number_of_edges, - rmm::cuda_stream_view stream_view) +void fill_offset( + VT* source, ET* offsets, VT number_of_vertices, ET number_of_edges, cuda::stream_ref stream_view) { cugraph::fill( rmm::exec_policy(stream_view), offsets, offsets + number_of_vertices + 1, number_of_edges); @@ -120,7 +118,7 @@ template rmm::device_buffer create_offset(VT* source, VT number_of_vertices, ET number_of_edges, - rmm::cuda_stream_view stream_view, + cuda::stream_ref stream_view, rmm::device_async_resource_ref mr) { // Offset array needs an extra element at the end to contain the ending offsets @@ -139,9 +137,9 @@ template std::unique_ptr> coo_to_csr( legacy::GraphCOOView const& graph, rmm::device_async_resource_ref mr) { - rmm::cuda_stream_view stream_view; + cuda::stream_ref stream_view; - legacy::GraphCOO temp_graph(graph, stream_view.value(), mr); + legacy::GraphCOO temp_graph(graph, stream_view.get(), mr); legacy::GraphCOOView temp_graph_view = temp_graph.view(); VT total_vertex_count = detail::sort(temp_graph_view, stream_view); rmm::device_buffer offsets = detail::create_offset( @@ -157,7 +155,7 @@ std::unique_ptr> coo_to_csr( // All conversion work runs on stream_view, which callers do not share (for example // raft::handle_t defaults to cudaStreamPerThread). Synchronize before returning so // downstream algorithms can safely consume the CSR buffers on any stream. - stream_view.synchronize(); + stream_view.sync(); return std::make_unique>(std::move(csr_contents)); } diff --git a/cpp/src/cores/core_number_impl.cuh b/cpp/src/cores/core_number_impl.cuh index 5d4d976a59..3f99fe073f 100644 --- a/cpp/src/cores/core_number_impl.cuh +++ b/cpp/src/cores/core_number_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/src/dag/topological_sort_impl.cuh b/cpp/src/dag/topological_sort_impl.cuh index 9f2164f0fb..5b0c15d2f4 100644 --- a/cpp/src/dag/topological_sort_impl.cuh +++ b/cpp/src/dag/topological_sort_impl.cuh @@ -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 */ diff --git a/cpp/src/detail/device_comm_wrapper.cuh b/cpp/src/detail/device_comm_wrapper.cuh index 97d31b9ffa..df6da4df87 100644 --- a/cpp/src/detail/device_comm_wrapper.cuh +++ b/cpp/src/detail/device_comm_wrapper.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 */ diff --git a/cpp/src/detail/permute_range.cuh b/cpp/src/detail/permute_range.cuh index 2cd1892a84..a166b31c3f 100644 --- a/cpp/src/detail/permute_range.cuh +++ b/cpp/src/detail/permute_range.cuh @@ -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 */ @@ -160,7 +160,7 @@ rmm::device_uvector permute_range(raft::handle_t const& handle, raft::copy(permuted_integers.data() + local_range_size - nr_deficits, extra_cluster_ids.begin() + deficits[comm_rank], nr_deficits, - handle.get_stream()); + handle.get_stream().get()); } assert(permuted_integers.size() == local_range_size); diff --git a/cpp/src/detail/utility_wrappers_32_common.cu b/cpp/src/detail/utility_wrappers_32_common.cu index ee11f1a963..e43ccbfb9a 100644 --- a/cpp/src/detail/utility_wrappers_32_common.cu +++ b/cpp/src/detail/utility_wrappers_32_common.cu @@ -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 */ @@ -17,6 +17,7 @@ #include #include #include +#include #include #include #include @@ -28,14 +29,14 @@ namespace cugraph { namespace detail { -template CUGRAPH_EXPORT void uniform_random_fill(rmm::cuda_stream_view const& stream_view, +template CUGRAPH_EXPORT void uniform_random_fill(cuda::stream_ref const& stream_view, int32_t* d_value, size_t size, int32_t min_value, int32_t max_value, raft::random::RngState& rng_state); -template CUGRAPH_EXPORT void uniform_random_fill(rmm::cuda_stream_view const& stream_view, +template CUGRAPH_EXPORT void uniform_random_fill(cuda::stream_ref const& stream_view, float* d_value, size_t size, float min_value, @@ -44,14 +45,14 @@ template CUGRAPH_EXPORT void uniform_random_fill(rmm::cuda_stream_view const& st template CUGRAPH_EXPORT void transform_increment_ints(raft::device_span values, int32_t value, - rmm::cuda_stream_view const& stream_view); + cuda::stream_ref const& stream_view); template CUGRAPH_EXPORT void transform_not_equal(raft::device_span values, raft::device_span result, int32_t compare, - rmm::cuda_stream_view const& stream_view); + cuda::stream_ref const& stream_view); -template CUGRAPH_EXPORT int32_t compute_maximum_vertex_id(rmm::cuda_stream_view const& stream_view, +template CUGRAPH_EXPORT int32_t compute_maximum_vertex_id(cuda::stream_ref const& stream_view, int32_t const* d_edgelist_srcs, int32_t const* d_edgelist_dsts, size_t num_edges); diff --git a/cpp/src/detail/utility_wrappers_64_common.cu b/cpp/src/detail/utility_wrappers_64_common.cu index 621d88ad9a..109a429dd7 100644 --- a/cpp/src/detail/utility_wrappers_64_common.cu +++ b/cpp/src/detail/utility_wrappers_64_common.cu @@ -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 */ @@ -17,6 +17,7 @@ #include #include #include +#include #include #include #include @@ -28,14 +29,14 @@ namespace cugraph { namespace detail { -template CUGRAPH_EXPORT void uniform_random_fill(rmm::cuda_stream_view const& stream_view, +template CUGRAPH_EXPORT void uniform_random_fill(cuda::stream_ref const& stream_view, int64_t* d_value, size_t size, int64_t min_value, int64_t max_value, raft::random::RngState& rng_state); -template CUGRAPH_EXPORT void uniform_random_fill(rmm::cuda_stream_view const& stream_view, +template CUGRAPH_EXPORT void uniform_random_fill(cuda::stream_ref const& stream_view, double* d_value, size_t size, double min_value, @@ -44,14 +45,14 @@ template CUGRAPH_EXPORT void uniform_random_fill(rmm::cuda_stream_view const& st template CUGRAPH_EXPORT void transform_increment_ints(raft::device_span values, int64_t value, - rmm::cuda_stream_view const& stream_view); + cuda::stream_ref const& stream_view); template CUGRAPH_EXPORT void transform_not_equal(raft::device_span values, raft::device_span result, int64_t compare, - rmm::cuda_stream_view const& stream_view); + cuda::stream_ref const& stream_view); -template CUGRAPH_EXPORT int64_t compute_maximum_vertex_id(rmm::cuda_stream_view const& stream_view, +template CUGRAPH_EXPORT int64_t compute_maximum_vertex_id(cuda::stream_ref const& stream_view, int64_t const* d_edgelist_srcs, int64_t const* d_edgelist_dsts, size_t num_edges); diff --git a/cpp/src/detail/utility_wrappers_impl.cuh b/cpp/src/detail/utility_wrappers_impl.cuh index 40c10fc001..7e6d6274be 100644 --- a/cpp/src/detail/utility_wrappers_impl.cuh +++ b/cpp/src/detail/utility_wrappers_impl.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 */ @@ -16,6 +16,7 @@ #include #include #include +#include #include #include #include @@ -29,7 +30,7 @@ namespace cugraph { namespace detail { template -void uniform_random_fill(rmm::cuda_stream_view const& stream_view, +void uniform_random_fill(cuda::stream_ref const& stream_view, value_t* d_value, size_t size, value_t min_value, @@ -38,17 +39,17 @@ void uniform_random_fill(rmm::cuda_stream_view const& stream_view, { if constexpr (std::is_integral::value) { raft::random::uniformInt( - rng_state, d_value, size, min_value, max_value, stream_view.value()); + rng_state, d_value, size, min_value, max_value, stream_view.get()); } else { raft::random::uniform( - rng_state, d_value, size, min_value, max_value, stream_view.value()); + rng_state, d_value, size, min_value, max_value, stream_view.get()); } } template void transform_increment_ints(raft::device_span values, value_t incr, - rmm::cuda_stream_view const& stream_view) + cuda::stream_ref const& stream_view) { thrust::transform(rmm::exec_policy(stream_view), values.begin(), @@ -63,7 +64,7 @@ template void transform_not_equal(raft::device_span values, raft::device_span result, value_t compare, - rmm::cuda_stream_view const& stream_view) + cuda::stream_ref const& stream_view) { thrust::transform(rmm::exec_policy(stream_view), values.begin(), @@ -74,7 +75,7 @@ void transform_not_equal(raft::device_span values, } template -vertex_t compute_maximum_vertex_id(rmm::cuda_stream_view const& stream_view, +vertex_t compute_maximum_vertex_id(cuda::stream_ref const& stream_view, vertex_t const* d_edgelist_srcs, vertex_t const* d_edgelist_dsts, size_t num_edges) diff --git a/cpp/src/layout/legacy/barnes_hut.cuh b/cpp/src/layout/legacy/barnes_hut.cuh index 77efaa1563..22c46e58e3 100644 --- a/cpp/src/layout/legacy/barnes_hut.cuh +++ b/cpp/src/layout/legacy/barnes_hut.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -20,6 +20,7 @@ #include #include +#include #include #include @@ -50,7 +51,7 @@ void barnes_hut(raft::handle_t const& handle, bool verbose = false, internals::GraphBasedDimRedCallback* callback = nullptr) { - rmm::cuda_stream_view stream_view(handle.get_stream()); + cuda::stream_ref stream_view(handle.get_stream()); const edge_t e = graph.number_of_edges; const vertex_t n = graph.number_of_vertices; @@ -76,8 +77,8 @@ void barnes_hut(raft::handle_t const& handle, int* bottomd = d_bottomd.data(); float* radiusd = d_radiusd.data(); - InitializationKernel<<<1, 1, 0, stream_view.value()>>>(limiter, maxdepthd, radiusd); - RAFT_CHECK_CUDA(stream_view.value()); + InitializationKernel<<<1, 1, 0, stream_view.get()>>>(limiter, maxdepthd, radiusd); + RAFT_CHECK_CUDA(stream_view.get()); const int FOUR_NNODES = 4 * nnodes; const int FOUR_N = 4 * n; @@ -124,8 +125,8 @@ void barnes_hut(raft::handle_t const& handle, // Copy start x and y positions. if (x_start && y_start) { - raft::copy(nodes_pos, x_start, n, stream_view.value()); - raft::copy(nodes_pos + nnodes + 1, y_start, n, stream_view.value()); + raft::copy(nodes_pos, x_start, n, stream_view.get()); + raft::copy(nodes_pos + nnodes + 1, y_start, n, stream_view.get()); } else { // Initialize positions with random values uniform_random_fill( @@ -152,21 +153,21 @@ void barnes_hut(raft::handle_t const& handle, if (graph.number_of_edges > 0) { // Sort COO for coalesced memory access. - sort(graph, stream_view.value()); - RAFT_CHECK_CUDA(stream_view.value()); + sort(graph, stream_view.get()); + RAFT_CHECK_CUDA(stream_view.get()); } if (vertex_mass != nullptr) { // Fill masses with 1 (because `nnodes + 1 > n`) cugraph::fill(handle.get_thrust_policy(), d_massl.begin() + n, d_massl.end(), 1.f); - raft::copy(massl, vertex_mass, n, stream_view.value()); + raft::copy(massl, vertex_mass, n, stream_view.get()); } else { // FA2 requires degree + 1 d_massl_edge_t.resize(nnodes + 1, handle.get_stream()); cugraph::fill(handle.get_thrust_policy(), d_massl_edge_t.begin(), d_massl_edge_t.end(), 1); massl_edge_t = d_massl_edge_t.data(); graph.degree(massl_edge_t, cugraph::legacy::DegreeDirection::OUT); - RAFT_CHECK_CUDA(stream_view.value()); + RAFT_CHECK_CUDA(stream_view.get()); thrust::transform( handle.get_thrust_policy(), @@ -216,66 +217,65 @@ void barnes_hut(raft::handle_t const& handle, cugraph::fill(handle.get_thrust_policy(), d_swinging.begin(), d_swinging.end(), 0.f); cugraph::fill(handle.get_thrust_policy(), d_traction.begin(), d_traction.end(), 0.f); - ResetKernel<<<1, 1, 0, stream_view.value()>>>(radiusd_squared, bottomd, NNODES, radiusd); - RAFT_CHECK_CUDA(stream_view.value()); + ResetKernel<<<1, 1, 0, stream_view.get()>>>(radiusd_squared, bottomd, NNODES, radiusd); + RAFT_CHECK_CUDA(stream_view.get()); // Compute bounding box arround all bodies - BoundingBoxKernel<<>>( - startl, - childl, - massl, - nodes_pos, - nodes_pos + nnodes + 1, - maxxl, - maxyl, - minxl, - minyl, - FOUR_NNODES, - NNODES, - n, - limiter, - radiusd); - RAFT_CHECK_CUDA(stream_view.value()); - - ClearKernel1<<>>(childl, FOUR_NNODES, FOUR_N); - RAFT_CHECK_CUDA(stream_view.value()); + BoundingBoxKernel<<>>(startl, + childl, + massl, + nodes_pos, + nodes_pos + nnodes + 1, + maxxl, + maxyl, + minxl, + minyl, + FOUR_NNODES, + NNODES, + n, + limiter, + radiusd); + RAFT_CHECK_CUDA(stream_view.get()); + + ClearKernel1<<>>(childl, FOUR_NNODES, FOUR_N); + RAFT_CHECK_CUDA(stream_view.get()); // Build quadtree - TreeBuildingKernel<<>>( + TreeBuildingKernel<<>>( childl, nodes_pos, nodes_pos + nnodes + 1, NNODES, n, maxdepthd, bottomd, radiusd); - RAFT_CHECK_CUDA(stream_view.value()); + RAFT_CHECK_CUDA(stream_view.get()); - ClearKernel2<<>>(startl, massl, NNODES, bottomd); - RAFT_CHECK_CUDA(stream_view.value()); + ClearKernel2<<>>(startl, massl, NNODES, bottomd); + RAFT_CHECK_CUDA(stream_view.get()); // Summarizes mass and position for each cell, bottom up approach - SummarizationKernel<<>>( + SummarizationKernel<<>>( countl, childl, massl, nodes_pos, nodes_pos + nnodes + 1, NNODES, n, bottomd); - RAFT_CHECK_CUDA(stream_view.value()); + RAFT_CHECK_CUDA(stream_view.get()); // Group closed bodies together, used to speed up Repulsion kernel - SortKernel<<>>( + SortKernel<<>>( sortl, countl, startl, childl, NNODES, n, bottomd); - RAFT_CHECK_CUDA(stream_view.value()); + RAFT_CHECK_CUDA(stream_view.get()); // Force computation O(n . log(n)) - RepulsionKernel<<>>(scaling_ratio, - theta, - epssq, - sortl, - childl, - massl, - nodes_pos, - nodes_pos + nnodes + 1, - rep_forces, - rep_forces + nnodes + 1, - theta_squared, - NNODES, - FOUR_NNODES, - n, - radiusd_squared, - maxdepthd); - RAFT_CHECK_CUDA(stream_view.value()); + RepulsionKernel<<>>(scaling_ratio, + theta, + epssq, + sortl, + childl, + massl, + nodes_pos, + nodes_pos + nnodes + 1, + rep_forces, + rep_forces + nnodes + 1, + theta_squared, + NNODES, + FOUR_NNODES, + n, + radiusd_squared, + maxdepthd); + RAFT_CHECK_CUDA(stream_view.get()); apply_gravity(nodes_pos, nodes_pos + nnodes + 1, @@ -286,7 +286,7 @@ void barnes_hut(raft::handle_t const& handle, strong_gravity_mode, scaling_ratio, n, - stream_view.value()); + stream_view.get()); apply_attraction(row, col, @@ -303,7 +303,7 @@ void barnes_hut(raft::handle_t const& handle, outbound_att_compensation, prevent_overlapping, vertex_radius, - stream_view.value()); + stream_view.get()); compute_local_speed(rep_forces, rep_forces + nnodes + 1, @@ -315,7 +315,7 @@ void barnes_hut(raft::handle_t const& handle, swinging, traction, n, - stream_view.value()); + stream_view.get()); // Compute global swinging and traction values const float s = @@ -328,19 +328,19 @@ void barnes_hut(raft::handle_t const& handle, adapt_speed(jitter_tolerance, &jt, &speed, &speed_efficiency, s, t, n); // Update positions - apply_forces_bh<<>>(nodes_pos, - nodes_pos + nnodes + 1, - attract, - attract + n, - rep_forces, - rep_forces + nnodes + 1, - old_forces, - old_forces + n, - swinging, - prevent_overlapping, - vertex_mobility, - speed, - n); + apply_forces_bh<<>>(nodes_pos, + nodes_pos + nnodes + 1, + attract, + attract + n, + rep_forces, + rep_forces + nnodes + 1, + old_forces, + old_forces + n, + swinging, + prevent_overlapping, + vertex_mobility, + speed, + n); if (callback) callback->on_epoch_end(nodes_pos); @@ -352,8 +352,8 @@ void barnes_hut(raft::handle_t const& handle, } // Copy nodes positions into final output pos - raft::copy(pos, nodes_pos, n, stream_view.value()); - raft::copy(pos + n, nodes_pos + nnodes + 1, n, stream_view.value()); + raft::copy(pos, nodes_pos, n, stream_view.get()); + raft::copy(pos + n, nodes_pos + nnodes + 1, n, stream_view.get()); if (callback) callback->on_train_end(nodes_pos); } diff --git a/cpp/src/layout/legacy/exact_fa2.cuh b/cpp/src/layout/legacy/exact_fa2.cuh index d0efabf170..a8f5613b3a 100644 --- a/cpp/src/layout/legacy/exact_fa2.cuh +++ b/cpp/src/layout/legacy/exact_fa2.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -81,25 +81,25 @@ void exact_fa2(raft::handle_t const& handle, uniform_random_fill(handle.get_stream(), pos, n * 2, -100.0f, 100.0f, rng_state); if (x_start && y_start) { - raft::copy(pos, x_start, n, stream_view.value()); - raft::copy(pos + n, y_start, n, stream_view.value()); + raft::copy(pos, x_start, n, stream_view.get()); + raft::copy(pos + n, y_start, n, stream_view.get()); } if (graph.number_of_edges > 0) { // Sort COO for coalesced memory access. - sort(graph, stream_view.value()); - RAFT_CHECK_CUDA(stream_view.value()); + sort(graph, stream_view.get()); + RAFT_CHECK_CUDA(stream_view.get()); } if (vertex_mass != nullptr) { - raft::copy(d_mass, vertex_mass, n, stream_view.value()); + raft::copy(d_mass, vertex_mass, n, stream_view.get()); } else { // FA2 requires degree + 1. mass_edge_t.resize(n, handle.get_stream()); cugraph::fill(handle.get_thrust_policy(), mass_edge_t.begin(), mass_edge_t.end(), 1); d_mass_edge_t = mass_edge_t.data(); graph.degree(d_mass_edge_t, cugraph::legacy::DegreeDirection::OUT); - RAFT_CHECK_CUDA(stream_view.value()); + RAFT_CHECK_CUDA(stream_view.get()); thrust::transform( handle.get_thrust_policy(), @@ -146,7 +146,7 @@ void exact_fa2(raft::handle_t const& handle, vertex_radius, overlap_scaling_ratio, n, - stream_view.value()); + stream_view.get()); apply_gravity(pos, pos + n, @@ -157,7 +157,7 @@ void exact_fa2(raft::handle_t const& handle, strong_gravity_mode, scaling_ratio, n, - stream_view.value()); + stream_view.get()); apply_attraction(row, col, @@ -174,7 +174,7 @@ void exact_fa2(raft::handle_t const& handle, outbound_att_compensation, prevent_overlapping, vertex_radius, - stream_view.value()); + stream_view.get()); compute_local_speed(d_repel, d_repel + n, @@ -186,7 +186,7 @@ void exact_fa2(raft::handle_t const& handle, d_swinging, d_traction, n, - stream_view.value()); + stream_view.get()); // Compute global swinging and traction values. const float s = thrust::reduce(handle.get_thrust_policy(), swinging.begin(), swinging.end()); @@ -207,7 +207,7 @@ void exact_fa2(raft::handle_t const& handle, vertex_mobility, speed, n, - stream_view.value()); + stream_view.get()); if (callback) callback->on_epoch_end(pos); diff --git a/cpp/src/linear_assignment/legacy/hungarian.cu b/cpp/src/linear_assignment/legacy/hungarian.cu index 2835d88e0f..6f0096f82c 100644 --- a/cpp/src/linear_assignment/legacy/hungarian.cu +++ b/cpp/src/linear_assignment/legacy/hungarian.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ // #define TIMING @@ -16,7 +16,6 @@ #include -#include #include #include diff --git a/cpp/src/link_analysis/pagerank_impl.cuh b/cpp/src/link_analysis/pagerank_impl.cuh index 2502c7fd81..ebbff656a3 100644 --- a/cpp/src/link_analysis/pagerank_impl.cuh +++ b/cpp/src/link_analysis/pagerank_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/src/link_prediction/detail/similarity_impl.cuh b/cpp/src/link_prediction/detail/similarity_impl.cuh index e59aada453..0187eae5af 100644 --- a/cpp/src/link_prediction/detail/similarity_impl.cuh +++ b/cpp/src/link_prediction/detail/similarity_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/src/lookup/lookup_src_dst_impl.cuh b/cpp/src/lookup/lookup_src_dst_impl.cuh index 39da5bae74..350b52e6ef 100644 --- a/cpp/src/lookup/lookup_src_dst_impl.cuh +++ b/cpp/src/lookup/lookup_src_dst_impl.cuh @@ -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 */ #include "detail/shuffle_wrappers.hpp" diff --git a/cpp/src/sampling/detail/check_edge_bias_values.cuh b/cpp/src/sampling/detail/check_edge_bias_values.cuh index 68bf946bfc..b5a2654fcc 100644 --- a/cpp/src/sampling/detail/check_edge_bias_values.cuh +++ b/cpp/src/sampling/detail/check_edge_bias_values.cuh @@ -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 */ diff --git a/cpp/src/sampling/negative_sampling_impl.cuh b/cpp/src/sampling/negative_sampling_impl.cuh index b89b108f5e..3350a6900c 100644 --- a/cpp/src/sampling/negative_sampling_impl.cuh +++ b/cpp/src/sampling/negative_sampling_impl.cuh @@ -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 */ @@ -548,12 +548,12 @@ std::tuple, rmm::device_uvector> negativ raft::copy(srcs.data() + num_samples - nr_deficits, extra_srcs.begin() + deficits[comm_rank], nr_deficits, - handle.get_stream()); + handle.get_stream().get()); raft::copy(dsts.data() + num_samples - nr_deficits, extra_dsts.begin() + deficits[comm_rank], nr_deficits, - handle.get_stream()); + handle.get_stream().get()); } return std::make_tuple(std::move(srcs), std::move(dsts)); diff --git a/cpp/src/sampling/random_walks_impl.cuh b/cpp/src/sampling/random_walks_impl.cuh index 2da96440bc..d3fc97a980 100644 --- a/cpp/src/sampling/random_walks_impl.cuh +++ b/cpp/src/sampling/random_walks_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -408,7 +408,7 @@ struct node2vec_selector { recv_counts = cugraph::host_scalar_allgather( handle.get_subcomm(cugraph::partition_manager::minor_comm_name()), - intersection_offsets.back_element(handle.get_stream()), + intersection_offsets.back_element(handle.get_stream().get()), handle.get_stream()); displacements.resize(recv_counts.size()); diff --git a/cpp/src/sampling/sampling_post_processing_impl.cuh b/cpp/src/sampling/sampling_post_processing_impl.cuh index 4e47a3e14e..b9446b3a84 100644 --- a/cpp/src/sampling/sampling_post_processing_impl.cuh +++ b/cpp/src/sampling/sampling_post_processing_impl.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 */ @@ -658,7 +658,7 @@ compute_min_hop_for_unique_label_vertex_pairs( h_label_offsets[i + 1] - h_label_offsets[i], offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); if (tmp_storage_bytes > d_tmp_storage.size()) { d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); @@ -672,7 +672,7 @@ compute_min_hop_for_unique_label_vertex_pairs( h_label_offsets[i + 1] - h_label_offsets[i], offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); } d_tmp_storage.resize(0, handle.get_stream()); d_tmp_storage.shrink_to_fit(handle.get_stream()); @@ -723,7 +723,7 @@ compute_min_hop_for_unique_label_vertex_pairs( (*seed_vertex_label_offsets).size() - 1, (*seed_vertex_label_offsets).begin(), (*seed_vertex_label_offsets).begin() + 1, - handle.get_stream()); + handle.get_stream().get()); if (tmp_storage_bytes > d_tmp_storage.size()) { d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); @@ -737,7 +737,7 @@ compute_min_hop_for_unique_label_vertex_pairs( (*seed_vertex_label_offsets).size() - 1, (*seed_vertex_label_offsets).begin(), (*seed_vertex_label_offsets).begin() + 1, - handle.get_stream()); + handle.get_stream().get()); /* enumerate unique (label, vertex) pairs */ @@ -1854,7 +1854,7 @@ renumber_sampled_edgelist(raft::handle_t const& handle, h_label_offsets[i + 1] - h_label_offsets[i], offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); if (tmp_storage_bytes > d_tmp_storage.size()) { d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); @@ -1870,7 +1870,7 @@ renumber_sampled_edgelist(raft::handle_t const& handle, h_label_offsets[i + 1] - h_label_offsets[i], offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); } new_vertices.resize(0, handle.get_stream()); d_tmp_storage.resize(0, handle.get_stream()); @@ -2126,7 +2126,7 @@ heterogeneous_renumber_sampled_edgelist( h_label_offsets[i + 1] - h_label_offsets[i], offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); if (tmp_storage_bytes > d_tmp_storage.size()) { d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); @@ -2143,7 +2143,7 @@ heterogeneous_renumber_sampled_edgelist( h_label_offsets[i + 1] - h_label_offsets[i], offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); } new_vertices.resize(0, handle.get_stream()); @@ -2340,7 +2340,7 @@ heterogeneous_renumber_sampled_edgelist( h_label_offsets[i + 1] - h_label_offsets[i], offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); if (tmp_storage_bytes > d_tmp_storage.size()) { d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); @@ -2357,7 +2357,7 @@ heterogeneous_renumber_sampled_edgelist( h_label_offsets[i + 1] - h_label_offsets[i], offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); } new_edge_ids.resize(0, handle.get_stream()); diff --git a/cpp/src/structure/coarsen_graph_impl.cuh b/cpp/src/structure/coarsen_graph_impl.cuh index ad6cfe60d1..5217a90f0a 100644 --- a/cpp/src/structure/coarsen_graph_impl.cuh +++ b/cpp/src/structure/coarsen_graph_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -28,6 +28,7 @@ #include #include #include +#include #include #include #include @@ -74,7 +75,7 @@ std::tuple, groupby_e_and_coarsen_edgelist(rmm::device_uvector&& edgelist_majors, rmm::device_uvector&& edgelist_minors, std::optional>&& edgelist_weights, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { auto pair_first = thrust::make_zip_iterator(edgelist_majors.begin(), edgelist_minors.begin()); diff --git a/cpp/src/structure/create_graph_from_edgelist_impl.cuh b/cpp/src/structure/create_graph_from_edgelist_impl.cuh index 22b8cef2f9..d3dafbb67f 100644 --- a/cpp/src/structure/create_graph_from_edgelist_impl.cuh +++ b/cpp/src/structure/create_graph_from_edgelist_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/src/structure/detail/structure_utils.cuh b/cpp/src/structure/detail/structure_utils.cuh index 7a86609e90..c9fbeafe03 100644 --- a/cpp/src/structure/detail/structure_utils.cuh +++ b/cpp/src/structure/detail/structure_utils.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -514,7 +514,7 @@ void sort_adjacency_list(raft::handle_t const& handle, h_vertex_offsets[i + 1] - h_vertex_offsets[i], offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); if (tmp_storage_bytes > d_tmp_storage.size()) { d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); } @@ -528,7 +528,7 @@ void sort_adjacency_list(raft::handle_t const& handle, h_vertex_offsets[i + 1] - h_vertex_offsets[i], offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); thrust::copy(handle.get_thrust_policy(), segment_sorted_indices.begin(), segment_sorted_indices.begin() + (h_edge_offsets[i + 1] - h_edge_offsets[i]), @@ -640,7 +640,7 @@ void sort_adjacency_list(raft::handle_t const& handle, h_vertex_offsets[i + 1] - h_vertex_offsets[i], offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); if (tmp_storage_bytes > d_tmp_storage.size()) { d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream()); } @@ -652,7 +652,7 @@ void sort_adjacency_list(raft::handle_t const& handle, h_vertex_offsets[i + 1] - h_vertex_offsets[i], offset_first, offset_first + 1, - handle.get_stream()); + handle.get_stream().get()); thrust::copy(handle.get_thrust_policy(), segment_sorted_indices.begin(), segment_sorted_indices.begin() + (h_edge_offsets[i + 1] - h_edge_offsets[i]), diff --git a/cpp/src/structure/edge_partition_device_view_impl.cuh b/cpp/src/structure/edge_partition_device_view_impl.cuh index 0d473d2a32..795a50407d 100644 --- a/cpp/src/structure/edge_partition_device_view_impl.cuh +++ b/cpp/src/structure/edge_partition_device_view_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -8,9 +8,9 @@ #include -#include #include +#include #include namespace cugraph { @@ -29,7 +29,7 @@ __host__ void compute_number_of_edges_with_mask_async_mg( vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { compute_number_of_edges_with_mask_async_mg(edge_mask, majors.begin(), @@ -51,7 +51,7 @@ __host__ void compute_number_of_edges_with_mask_async_mg( vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { compute_number_of_edges_with_mask_async_mg( edge_mask, @@ -74,7 +74,7 @@ __host__ void compute_number_of_edges_with_mask_async_mg( vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { auto major_first = cuda::make_transform_iterator(majors.offsets.data(), shift_right_t{majors.base_major}); @@ -97,7 +97,7 @@ __host__ rmm::device_uvector compute_local_degrees_with_mask_mg( vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { return compute_local_degrees_with_mask_mg(edge_mask, majors.begin(), @@ -117,7 +117,7 @@ __host__ rmm::device_uvector compute_local_degrees_with_mask_mg( vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { return compute_local_degrees_with_mask_mg( edge_mask, @@ -140,7 +140,7 @@ __host__ void compute_number_of_edges_with_mask_async_sg( raft::device_span majors, raft::device_span count, raft::device_span offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { compute_number_of_edges_with_mask_async_sg( edge_mask, majors.begin(), majors.end(), count, offsets, stream); @@ -152,7 +152,7 @@ __host__ void compute_number_of_edges_with_mask_async_sg( std::tuple vertex_partition_range, raft::device_span count, raft::device_span offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { compute_number_of_edges_with_mask_async_sg( edge_mask, @@ -168,7 +168,7 @@ __host__ rmm::device_uvector compute_local_degrees_with_mask_sg( cuda::std::optional edge_mask, raft::device_span majors, raft::device_span offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { return compute_local_degrees_with_mask_sg( edge_mask, majors.begin(), majors.end(), offsets, stream); @@ -179,7 +179,7 @@ __host__ rmm::device_uvector compute_local_degrees_with_mask_sg( cuda::std::optional edge_mask, std::tuple vertex_partition_range, raft::device_span offsets, - rmm::cuda_stream_view stream) + cuda::stream_ref stream) { return compute_local_degrees_with_mask_sg( edge_mask, diff --git a/cpp/src/structure/edge_partition_device_view_mask_degrees_common_v32_e32.cu b/cpp/src/structure/edge_partition_device_view_mask_degrees_common_v32_e32.cu index d931e1040e..b1421f95d5 100644 --- a/cpp/src/structure/edge_partition_device_view_mask_degrees_common_v32_e32.cu +++ b/cpp/src/structure/edge_partition_device_view_mask_degrees_common_v32_e32.cu @@ -15,6 +15,8 @@ #include #include +#include + namespace cugraph { namespace detail { @@ -25,7 +27,7 @@ template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degre cuda::std::optional edge_mask, raft::device_span majors, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degrees_with_mask_mg( cuda::std::optional edge_mask, @@ -34,14 +36,14 @@ template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degre vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_sg( cuda::std::optional edge_mask, raft::device_span majors, raft::device_span count, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg( cuda::std::optional edge_mask, @@ -51,7 +53,7 @@ template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); } // namespace detail } // namespace cugraph diff --git a/cpp/src/structure/edge_partition_device_view_mask_degrees_common_v64_e64.cu b/cpp/src/structure/edge_partition_device_view_mask_degrees_common_v64_e64.cu index 212c9d5ba4..18ba134757 100644 --- a/cpp/src/structure/edge_partition_device_view_mask_degrees_common_v64_e64.cu +++ b/cpp/src/structure/edge_partition_device_view_mask_degrees_common_v64_e64.cu @@ -15,6 +15,8 @@ #include #include +#include + namespace cugraph { namespace detail { @@ -25,7 +27,7 @@ template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degre cuda::std::optional edge_mask, raft::device_span majors, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degrees_with_mask_mg( cuda::std::optional edge_mask, @@ -34,14 +36,14 @@ template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degre vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_sg( cuda::std::optional edge_mask, raft::device_span majors, raft::device_span count, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg( cuda::std::optional edge_mask, @@ -51,7 +53,7 @@ template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); } // namespace detail } // namespace cugraph diff --git a/cpp/src/structure/edge_partition_device_view_mg_v32_e32.cu b/cpp/src/structure/edge_partition_device_view_mg_v32_e32.cu index 43a980a03e..faba5548fc 100644 --- a/cpp/src/structure/edge_partition_device_view_mg_v32_e32.cu +++ b/cpp/src/structure/edge_partition_device_view_mg_v32_e32.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include "edge_partition_device_view_impl.cuh" @@ -8,6 +8,8 @@ #include #include +#include + namespace cugraph { namespace detail { @@ -21,7 +23,7 @@ template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degre vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degrees_with_mask_mg( cuda::std::optional edge_mask, @@ -30,7 +32,7 @@ template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degre vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg( cuda::std::optional edge_mask, @@ -40,7 +42,7 @@ template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg( cuda::std::optional edge_mask, @@ -50,7 +52,7 @@ template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg( cuda::std::optional edge_mask, @@ -60,7 +62,7 @@ template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); } // namespace detail } // namespace cugraph diff --git a/cpp/src/structure/edge_partition_device_view_mg_v64_e64.cu b/cpp/src/structure/edge_partition_device_view_mg_v64_e64.cu index efb6bfa435..f75b539cc3 100644 --- a/cpp/src/structure/edge_partition_device_view_mg_v64_e64.cu +++ b/cpp/src/structure/edge_partition_device_view_mg_v64_e64.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include "edge_partition_device_view_impl.cuh" @@ -8,6 +8,8 @@ #include #include +#include + namespace cugraph { namespace detail { @@ -22,7 +24,7 @@ template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg( cuda::std::optional edge_mask, @@ -32,7 +34,7 @@ template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degrees_with_mask_mg( cuda::std::optional edge_mask, @@ -41,7 +43,7 @@ template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degre vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degrees_with_mask_mg( cuda::std::optional edge_mask, @@ -50,7 +52,7 @@ template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degre vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg( cuda::std::optional edge_mask, @@ -60,7 +62,7 @@ template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_mg vertex_t major_range_first, cuda::std::optional major_hypersparse_first, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); } // namespace detail } // namespace cugraph diff --git a/cpp/src/structure/edge_partition_device_view_sg_v32_e32.cu b/cpp/src/structure/edge_partition_device_view_sg_v32_e32.cu index 8046e4b14f..5524323358 100644 --- a/cpp/src/structure/edge_partition_device_view_sg_v32_e32.cu +++ b/cpp/src/structure/edge_partition_device_view_sg_v32_e32.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include "edge_partition_device_view_impl.cuh" @@ -8,6 +8,8 @@ #include #include +#include + namespace cugraph { namespace detail { @@ -19,26 +21,26 @@ template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_sg raft::device_span majors, raft::device_span count, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_sg( cuda::std::optional edge_mask, std::tuple vertex_partition_range, raft::device_span count, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degrees_with_mask_sg( cuda::std::optional edge_mask, raft::device_span majors, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degrees_with_mask_sg( cuda::std::optional edge_mask, std::tuple vertex_partition_range, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); } // namespace detail } // namespace cugraph diff --git a/cpp/src/structure/edge_partition_device_view_sg_v64_e64.cu b/cpp/src/structure/edge_partition_device_view_sg_v64_e64.cu index 9f8e41d38d..a1c51a0ba0 100644 --- a/cpp/src/structure/edge_partition_device_view_sg_v64_e64.cu +++ b/cpp/src/structure/edge_partition_device_view_sg_v64_e64.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include "edge_partition_device_view_impl.cuh" @@ -8,6 +8,8 @@ #include #include +#include + namespace cugraph { namespace detail { @@ -19,26 +21,26 @@ template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_sg raft::device_span majors, raft::device_span count, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ void compute_number_of_edges_with_mask_async_sg( cuda::std::optional edge_mask, std::tuple vertex_partition_range, raft::device_span count, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degrees_with_mask_sg( cuda::std::optional edge_mask, raft::device_span majors, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); template CUGRAPH_EXPORT __host__ rmm::device_uvector compute_local_degrees_with_mask_sg( cuda::std::optional edge_mask, std::tuple vertex_partition_range, raft::device_span offsets, - rmm::cuda_stream_view stream); + cuda::stream_ref stream); } // namespace detail } // namespace cugraph diff --git a/cpp/src/structure/graph_impl.cuh b/cpp/src/structure/graph_impl.cuh index 8a6aba81e9..fe6b2468e5 100644 --- a/cpp/src/structure/graph_impl.cuh +++ b/cpp/src/structure/graph_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/src/structure/graph_view_impl.cuh b/cpp/src/structure/graph_view_impl.cuh index 3df69a48c1..de0248c2a8 100644 --- a/cpp/src/structure/graph_view_impl.cuh +++ b/cpp/src/structure/graph_view_impl.cuh @@ -171,7 +171,7 @@ rmm::device_uvector compute_major_degrees( static_cast(num_local_degrees), raft::comms::op_t::SUM, i, - handle.get_stream()); + handle.get_stream().get()); } return degrees; @@ -328,7 +328,7 @@ edge_t count_edge_partition_multi_edges( cugraph::for_all_major_for_all_nbr_high_degree<<>>( + handle.get_stream().get()>>>( edge_partition, edge_partition.major_range_first(), edge_partition.major_range_first() + (*segment_offsets)[1], @@ -342,7 +342,7 @@ edge_t count_edge_partition_multi_edges( cugraph::for_all_major_for_all_nbr_mid_degree<<>>( + handle.get_stream().get()>>>( edge_partition, edge_partition.major_range_first() + (*segment_offsets)[1], edge_partition.major_range_first() + (*segment_offsets)[2], diff --git a/cpp/src/structure/graph_weight_utils_impl.cuh b/cpp/src/structure/graph_weight_utils_impl.cuh index 6caa5bb838..e362ab8477 100644 --- a/cpp/src/structure/graph_weight_utils_impl.cuh +++ b/cpp/src/structure/graph_weight_utils_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/src/structure/legacy/graph.cu b/cpp/src/structure/legacy/graph.cu index d781ee25b5..c9fc8d2806 100644 --- a/cpp/src/structure/legacy/graph.cu +++ b/cpp/src/structure/legacy/graph.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -15,6 +15,7 @@ #include #include +#include #include #include #include @@ -26,7 +27,7 @@ template void degree_from_offsets(vertex_t number_of_vertices, edge_t const* offsets, edge_t* degree, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { // Computes out-degree for x = 0 and x = 2 thrust::for_each( @@ -42,7 +43,7 @@ void degree_from_vertex_ids(const raft::handle_t* handle, edge_t number_of_edges, vertex_t const* indices, edge_t* degree, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { thrust::for_each( rmm::exec_policy(stream_view), @@ -51,7 +52,7 @@ void degree_from_vertex_ids(const raft::handle_t* handle, [indices, degree] __device__(edge_t e) { atomicAdd(degree + indices[e], edge_t{1}); }); if ((handle != nullptr) && (handle->comms_initialized())) { auto& comm = handle->get_comms(); - comm.allreduce(degree, degree, number_of_vertices, raft::comms::op_t::SUM, stream_view.value()); + comm.allreduce(degree, degree, number_of_vertices, raft::comms::op_t::SUM, stream_view.get()); } } @@ -75,7 +76,7 @@ template void GraphCompressedSparseBaseView::get_source_indices(VT* src_indices) const { CUGRAPH_EXPECTS(offsets != nullptr, "No graph specified"); - rmm::cuda_stream_view stream_view; + cuda::stream_ref stream_view; raft::device_span indices_span(src_indices, GraphViewBase::number_of_edges); @@ -145,7 +146,7 @@ void GraphCompressedSparseBaseView::degree(ET* degree, DegreeDirecti // (e.g. if you have a CSC and you want in-degree (x=1) then pass // the offsets/indices and request an out-degree (x=2)) // - rmm::cuda_stream_view stream_view; + cuda::stream_ref stream_view; if (direction != DegreeDirection::IN) { if ((GraphViewBase::handle != nullptr) && diff --git a/cpp/src/structure/renumber_edgelist_impl.cuh b/cpp/src/structure/renumber_edgelist_impl.cuh index 81db2eeedb..e2eb2a39f4 100644 --- a/cpp/src/structure/renumber_edgelist_impl.cuh +++ b/cpp/src/structure/renumber_edgelist_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -640,23 +640,23 @@ compute_renumber_map(raft::handle_t const& handle, auto edge_partition_major_range_sizes = host_scalar_allgather(minor_comm, sorted_local_vertices.size(), handle.get_stream()); for (int i = 0; i < minor_comm_size; ++i) { - auto sorted_majors = - large_vertex_buffer_type - ? large_buffer_manager::allocate_memory_buffer( - edge_partition_major_range_sizes[i], handle.get_stream()) - : rmm::device_uvector(edge_partition_major_range_sizes[i], handle.get_stream()); + auto sorted_majors = large_vertex_buffer_type + ? large_buffer_manager::allocate_memory_buffer( + edge_partition_major_range_sizes[i], handle.get_stream().get()) + : rmm::device_uvector(edge_partition_major_range_sizes[i], + handle.get_stream().get()); device_bcast(minor_comm, sorted_local_vertices.data(), sorted_majors.data(), edge_partition_major_range_sizes[i], i, - handle.get_stream()); + handle.get_stream().get()); auto sorted_major_degrees = large_vertex_buffer_type ? large_buffer_manager::allocate_memory_buffer(sorted_majors.size(), - handle.get_stream()) - : rmm::device_uvector(sorted_majors.size(), handle.get_stream()); + handle.get_stream().get()) + : rmm::device_uvector(sorted_majors.size(), handle.get_stream().get()); if (large_vertex_buffer_type) { compute_sorted_local_major_degrees_without_atomics( handle, @@ -691,7 +691,7 @@ compute_renumber_map(raft::handle_t const& handle, edge_partition_major_range_sizes[i], raft::comms::op_t::SUM, i, - handle.get_stream()); + handle.get_stream().get()); if (i == minor_comm_rank) { sorted_local_vertex_degrees = std::move(sorted_major_degrees); } } } else { @@ -1025,8 +1025,10 @@ std::vector aggregate_offset_vectors(raft::handle_t const& handle, raft::update_device(d_offsets.data(), offsets.data(), offsets.size(), handle.get_stream()); rmm::device_uvector d_aggregate_offset_vectors(minor_comm_size * d_offsets.size(), handle.get_stream()); - minor_comm.allgather( - d_offsets.data(), d_aggregate_offset_vectors.data(), d_offsets.size(), handle.get_stream()); + minor_comm.allgather(d_offsets.data(), + d_aggregate_offset_vectors.data(), + d_offsets.size(), + handle.get_stream().get()); std::vector h_aggregate_offset_vectors(d_aggregate_offset_vectors.size(), vertex_t{0}); raft::update_host(h_aggregate_offset_vectors.data(), diff --git a/cpp/src/structure/renumber_utils_impl.cuh b/cpp/src/structure/renumber_utils_impl.cuh index 992bd5abf9..46d3d9a874 100644 --- a/cpp/src/structure/renumber_utils_impl.cuh +++ b/cpp/src/structure/renumber_utils_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/src/structure/select_random_vertices_impl.cuh b/cpp/src/structure/select_random_vertices_impl.cuh index ce8d07acc7..02cd9db70a 100644 --- a/cpp/src/structure/select_random_vertices_impl.cuh +++ b/cpp/src/structure/select_random_vertices_impl.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 */ #pragma once diff --git a/cpp/src/traversal/bfs_impl.cuh b/cpp/src/traversal/bfs_impl.cuh index 2ae816ab8c..926fe3a6ef 100644 --- a/cpp/src/traversal/bfs_impl.cuh +++ b/cpp/src/traversal/bfs_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/src/traversal/extract_bfs_paths_impl.cuh b/cpp/src/traversal/extract_bfs_paths_impl.cuh index be7f395e84..b9e997ccf0 100644 --- a/cpp/src/traversal/extract_bfs_paths_impl.cuh +++ b/cpp/src/traversal/extract_bfs_paths_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/src/traversal/k_hop_nbrs_impl.cuh b/cpp/src/traversal/k_hop_nbrs_impl.cuh index c6232464e1..fd18d3fb9b 100644 --- a/cpp/src/traversal/k_hop_nbrs_impl.cuh +++ b/cpp/src/traversal/k_hop_nbrs_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/src/traversal/od_shortest_distances_impl.cuh b/cpp/src/traversal/od_shortest_distances_impl.cuh index a03e102590..e8dd717167 100644 --- a/cpp/src/traversal/od_shortest_distances_impl.cuh +++ b/cpp/src/traversal/od_shortest_distances_impl.cuh @@ -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 */ #pragma once @@ -813,7 +813,7 @@ rmm::device_uvector od_shortest_distances( multi_partition_copy_block_size, handle.get_device_properties().maxGridSize[0]); multi_partition_copy(1 /* near queue */ + num_far_buffers)> - <<>>( + <<>>( input_first + num_copied, input_first + num_copied + this_loop_size, raft::device_span(d_buffer_ptrs.data(), d_buffer_ptrs.size()), @@ -951,28 +951,30 @@ rmm::device_uvector od_shortest_distances( handle.get_device_properties().maxGridSize[0]); auto constexpr max_num_partitions = static_cast(1 /* near queue */ + num_far_buffers); - multi_partition_copy - <<>>( - tmp_buffer.begin(), - tmp_buffer.end(), - raft::device_span(d_buffer_ptrs.data(), d_buffer_ptrs.size()), - [key_to_dist_map = - detail::kv_cuco_store_find_device_view_t(key_to_dist_map.view()), - split_thresholds = raft::device_span( - d_split_thresholds.data(), d_split_thresholds.size()), - invalid_threshold] __device__(auto key) { - auto dist = key_to_dist_map.find(key); - return static_cast( - (dist < invalid_threshold) - ? max_num_partitions /* discard */ - : cuda::std::distance(split_thresholds.begin(), - thrust::upper_bound(thrust::seq, - split_thresholds.begin(), - split_thresholds.end(), - dist))); - }, - cuda::std::identity{}, - raft::device_span(d_counters.data(), d_counters.size())); + multi_partition_copy<<>>( + tmp_buffer.begin(), + tmp_buffer.end(), + raft::device_span(d_buffer_ptrs.data(), d_buffer_ptrs.size()), + [key_to_dist_map = + detail::kv_cuco_store_find_device_view_t(key_to_dist_map.view()), + split_thresholds = raft::device_span(d_split_thresholds.data(), + d_split_thresholds.size()), + invalid_threshold] __device__(auto key) { + auto dist = key_to_dist_map.find(key); + return static_cast( + (dist < invalid_threshold) + ? max_num_partitions /* discard */ + : cuda::std::distance(split_thresholds.begin(), + thrust::upper_bound(thrust::seq, + split_thresholds.begin(), + split_thresholds.end(), + dist))); + }, + cuda::std::identity{}, + raft::device_span(d_counters.data(), d_counters.size())); } std::vector h_counters(d_counters.size()); raft::update_host( diff --git a/cpp/src/traversal/sssp_impl.cuh b/cpp/src/traversal/sssp_impl.cuh index 1087f0afc8..41f3d1ce08 100644 --- a/cpp/src/traversal/sssp_impl.cuh +++ b/cpp/src/traversal/sssp_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once diff --git a/cpp/src/tree/legacy/mst.cu b/cpp/src/tree/legacy/mst.cu index b432f7cca2..62bf00b3d9 100644 --- a/cpp/src/tree/legacy/mst.cu +++ b/cpp/src/tree/legacy/mst.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -30,7 +30,7 @@ std::unique_ptr> mst_impl( rmm::device_async_resource_ref mr) { - auto stream = handle.get_stream(); + auto stream = handle.get_stream().get(); rmm::device_uvector colors(graph.number_of_vertices, stream); auto mst_edges = raft::sparse::solver::mst(handle, graph.offsets, diff --git a/cpp/src/utilities/partition_scatter_map_wrappers.cu b/cpp/src/utilities/partition_scatter_map_wrappers.cu index 60420546d1..b67e57ed3d 100644 --- a/cpp/src/utilities/partition_scatter_map_wrappers.cu +++ b/cpp/src/utilities/partition_scatter_map_wrappers.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 * * Explicit instantiations for cugraph/utilities/partition_scatter_map_wrappers.cuh. @@ -13,6 +13,7 @@ #include #include +#include #include #include @@ -28,7 +29,7 @@ rmm::device_uvector compute_partition_scatter_map_impl( offset_t const* intra_partition_displs, size_t const* group_displacements, size_t num_elements, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { rmm::device_uvector scatter_map(num_elements, stream_view); thrust::transform( @@ -51,7 +52,7 @@ rmm::device_uvector compute_partition_scatter_map_impl( offset_t const* intra_partition_displs, \ size_t const* group_displacements, \ size_t num_elements, \ - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) CUGRAPH_PARTITION_SCATTER_MAP_INST(std::uint8_t, std::uint32_t); CUGRAPH_PARTITION_SCATTER_MAP_INST(int, std::size_t); diff --git a/cpp/src/utilities/permute_wrappers.cu b/cpp/src/utilities/permute_wrappers.cu index 436ae7edf6..3784aa98ec 100644 --- a/cpp/src/utilities/permute_wrappers.cu +++ b/cpp/src/utilities/permute_wrappers.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 * * Explicit instantiations for cugraph/utilities/permute_wrappers.cuh. @@ -12,6 +12,7 @@ #include #include +#include #include #include @@ -24,7 +25,7 @@ template void permute_in_place_impl(T* first, std::size_t const* map_first, std::size_t num_elements, - rmm::cuda_stream_view stream_view) + cuda::stream_ref stream_view) { auto const policy = rmm::exec_policy(stream_view); rmm::device_uvector tmp(num_elements, stream_view); @@ -32,12 +33,11 @@ void permute_in_place_impl(T* first, thrust::copy(policy, tmp.begin(), tmp.end(), first); } -#define CUGRAPH_PERMUTE_IN_PLACE_SCALAR_INST(ScalarType) \ - template CUGRAPH_EXPORT void permute_in_place_impl( \ - ScalarType * first, \ - std::size_t const* map_first, \ - std::size_t num_elements, \ - rmm::cuda_stream_view stream_view) +#define CUGRAPH_PERMUTE_IN_PLACE_SCALAR_INST(ScalarType) \ + template CUGRAPH_EXPORT void permute_in_place_impl(ScalarType * first, \ + std::size_t const* map_first, \ + std::size_t num_elements, \ + cuda::stream_ref stream_view) CUGRAPH_PERMUTE_IN_PLACE_SCALAR_INST(std::int32_t); CUGRAPH_PERMUTE_IN_PLACE_SCALAR_INST(std::int64_t); diff --git a/cpp/src/utilities/shuffle_vertex_pairs.cuh b/cpp/src/utilities/shuffle_vertex_pairs.cuh index c437a30112..5f97309916 100644 --- a/cpp/src/utilities/shuffle_vertex_pairs.cuh +++ b/cpp/src/utilities/shuffle_vertex_pairs.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 */ diff --git a/cpp/src/utilities/validation_checks_impl.cuh b/cpp/src/utilities/validation_checks_impl.cuh index 7461258fe6..b4fa1bbb2e 100644 --- a/cpp/src/utilities/validation_checks_impl.cuh +++ b/cpp/src/utilities/validation_checks_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/tests/c_api/mg_test_utils.cpp b/cpp/tests/c_api/mg_test_utils.cpp index f94ace13f1..f96bd8fa3d 100644 --- a/cpp/tests/c_api/mg_test_utils.cpp +++ b/cpp/tests/c_api/mg_test_utils.cpp @@ -25,6 +25,8 @@ #include #include +#include + #include #include @@ -121,8 +123,9 @@ extern "C" void* create_mg_raft_handle(int argc, char** argv) // Match initialize_mg_handle: per-thread default stream + stream pool. constexpr size_t stream_pool_size = 8; // default CUDA_DEVICE_MAX_CONNECTIONS - raft::handle_t* handle = new raft::handle_t{ - rmm::cuda_stream_per_thread, std::make_shared(stream_pool_size)}; + raft::handle_t* handle = + new raft::handle_t{cuda::stream_ref{cudaStreamPerThread}, + std::make_shared(stream_pool_size)}; raft::comms::initialize_mpi_comms(handle, MPI_COMM_WORLD); #if 1 diff --git a/cpp/tests/community/egonet_test.cpp b/cpp/tests/community/egonet_test.cpp index 6dcc8f7f7c..5dd3c3ecde 100644 --- a/cpp/tests/community/egonet_test.cpp +++ b/cpp/tests/community/egonet_test.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include "community/egonet_validate.hpp" @@ -18,6 +18,8 @@ #include #include +#include + #include struct Egonet_Usecase { @@ -47,7 +49,7 @@ class Tests_Egonet : public ::testing::TestWithParam(n_streams); - raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); + raft::handle_t handle(cuda::stream_ref{cudaStreamPerThread}, stream_pool); bool renumber = true; diff --git a/cpp/tests/prims/mg_per_v_transform_reduce_dst_key_aggregated_outgoing_e.cu b/cpp/tests/prims/mg_per_v_transform_reduce_dst_key_aggregated_outgoing_e.cu index c8344afbbe..aab408ff49 100644 --- a/cpp/tests/prims/mg_per_v_transform_reduce_dst_key_aggregated_outgoing_e.cu +++ b/cpp/tests/prims/mg_per_v_transform_reduce_dst_key_aggregated_outgoing_e.cu @@ -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 */ @@ -34,6 +34,7 @@ #include #include #include +#include #include #include #include @@ -174,7 +175,7 @@ class Tests_MGPerVTransformReduceDstKeyAggregatedOutgoingE std::array reduction_types = { reduction_type_t::PLUS, reduction_type_t::ELEMWISE_MIN, reduction_type_t::ELEMWISE_MAX}; - std::vector(0, rmm::cuda_stream_view{}))> + std::vector(0, cuda::stream_ref{}))> mg_results{}; mg_results.reserve(reduction_types.size()); diff --git a/cpp/tests/prims/mg_per_v_transform_reduce_incoming_outgoing_e.cu b/cpp/tests/prims/mg_per_v_transform_reduce_incoming_outgoing_e.cu index b71cfe453a..57e1fd8c27 100644 --- a/cpp/tests/prims/mg_per_v_transform_reduce_incoming_outgoing_e.cu +++ b/cpp/tests/prims/mg_per_v_transform_reduce_incoming_outgoing_e.cu @@ -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 */ @@ -32,6 +32,7 @@ #include #include #include +#include #include #include #include @@ -141,9 +142,9 @@ class Tests_MGPerVTransformReduceIncomingOutgoingE std::array reduction_types = { reduction_type_t::PLUS, reduction_type_t::ELEMWISE_MIN, reduction_type_t::ELEMWISE_MAX}; - std::vector(0, rmm::cuda_stream_view{}))> + std::vector(0, cuda::stream_ref{}))> mg_in_results{}; - std::vector(0, rmm::cuda_stream_view{}))> + std::vector(0, cuda::stream_ref{}))> mg_out_results{}; mg_in_results.reserve(reduction_types.size()); mg_out_results.reserve(reduction_types.size()); diff --git a/cpp/tests/structure/streams.cu b/cpp/tests/structure/streams.cu index e28b278889..d81bbb588c 100644 --- a/cpp/tests/structure/streams.cu +++ b/cpp/tests/structure/streams.cu @@ -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 */ @@ -12,6 +12,7 @@ #include #include +#include #include struct StreamTest : public ::testing::Test {}; @@ -20,7 +21,7 @@ TEST_F(StreamTest, basic_test) { size_t n_streams = 4; auto stream_pool = std::make_shared(n_streams); - raft::handle_t handle(rmm::cuda_stream_per_thread, stream_pool); + raft::handle_t handle(cuda::stream_ref{cudaStreamPerThread}, stream_pool); const size_t input_size = 4096; @@ -37,7 +38,7 @@ TEST_F(StreamTest, basic_test) v.begin(), v.begin(), 2 * thrust::placeholders::_1 + thrust::placeholders::_2); - RAFT_CUDA_TRY(cudaStreamSynchronize(handle.get_next_usable_stream(i))); + handle.get_next_usable_stream(i).sync(); }, i); } diff --git a/cpp/tests/utilities/mg_utilities.cpp b/cpp/tests/utilities/mg_utilities.cpp index bd413a7c18..8e7c00434d 100644 --- a/cpp/tests/utilities/mg_utilities.cpp +++ b/cpp/tests/utilities/mg_utilities.cpp @@ -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 */ #include "utilities/mg_utilities.hpp" @@ -11,6 +11,8 @@ #include #include +#include + #include namespace cugraph { @@ -38,7 +40,7 @@ std::unique_ptr initialize_mg_handle(size_t pool_size) { std::unique_ptr handle{nullptr}; - handle = std::make_unique(rmm::cuda_stream_per_thread, + handle = std::make_unique(cuda::stream_ref{cudaStreamPerThread}, std::make_shared(pool_size)); auto comm_size = query_mpi_comm_world_size(); @@ -54,7 +56,7 @@ std::unique_ptr initialize_mg_handle(size_t pool_size) return std::move(handle); } -void enforce_p2p_initialization(raft::comms::comms_t const& comm, rmm::cuda_stream_view stream) +void enforce_p2p_initialization(raft::comms::comms_t const& comm, cuda::stream_ref stream) { auto const comm_size = comm.get_size(); diff --git a/cpp/tests/utilities/mg_utilities.hpp b/cpp/tests/utilities/mg_utilities.hpp index 63fe4a8390..2197fb7147 100644 --- a/cpp/tests/utilities/mg_utilities.hpp +++ b/cpp/tests/utilities/mg_utilities.hpp @@ -1,11 +1,13 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once #include +#include + #include namespace cugraph { @@ -23,7 +25,7 @@ std::unique_ptr initialize_mg_handle( // NCCL lazily initializes for P2P, and this enforces P2P initialization for better performance // measurements -void enforce_p2p_initialization(raft::comms::comms_t const& comm, rmm::cuda_stream_view stream); +void enforce_p2p_initialization(raft::comms::comms_t const& comm, cuda::stream_ref stream); } // namespace test } // namespace cugraph diff --git a/cpp/tests/utilities/property_generator_utilities.hpp b/cpp/tests/utilities/property_generator_utilities.hpp index f6f51fa88b..244166f81b 100644 --- a/cpp/tests/utilities/property_generator_utilities.hpp +++ b/cpp/tests/utilities/property_generator_utilities.hpp @@ -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 */ @@ -14,6 +14,8 @@ #include +#include + #include #include @@ -27,8 +29,8 @@ struct generate { using edge_type = typename GraphViewType::edge_type; using edge_type_t = int32_t; - using property_buffer_type = std::decay_t( - size_t{0}, rmm::cuda_stream_view{}))>; + using property_buffer_type = + std::decay_t(size_t{0}, cuda::stream_ref{}))>; public: static property_t initial_value(int32_t init); diff --git a/cpp/tests/utilities/test_graphs.hpp b/cpp/tests/utilities/test_graphs.hpp index f62438a4cb..ebe4054b29 100644 --- a/cpp/tests/utilities/test_graphs.hpp +++ b/cpp/tests/utilities/test_graphs.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once