Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Use raft::host_span instead of const reference to std::vector #4931

Merged
merged 18 commits into from
Feb 21, 2025
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
18 commits
Select commit Hold shift + click to select a range
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
151 changes: 76 additions & 75 deletions cpp/include/cugraph/utilities/device_comm.hpp

Large diffs are not rendered by default.

254 changes: 129 additions & 125 deletions cpp/include/cugraph/utilities/shuffle_comm.cuh

Large diffs are not rendered by default.

15 changes: 8 additions & 7 deletions cpp/src/c_api/extract_ego.cpp
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2022-2024, NVIDIA CORPORATION.
* Copyright (c) 2022-2025, NVIDIA CORPORATION.
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
Expand Down Expand Up @@ -154,12 +154,13 @@ struct extract_ego_functor : public cugraph::c_api::abstract_functor {
std::exclusive_scan(recvcounts.begin(), recvcounts.end(), displacements.begin(), size_t{0});
rmm::device_uvector<size_t> allgathered_indices(displacements.back() + recvcounts.back(),
handle_.get_stream());
cugraph::device_allgatherv(handle_.get_comms(),
(*source_indices).begin(),
allgathered_indices.begin(),
recvcounts,
displacements,
handle_.get_stream());
cugraph::device_allgatherv(
handle_.get_comms(),
(*source_indices).begin(),
allgathered_indices.begin(),
raft::host_span<size_t const>(recvcounts.data(), recvcounts.size()),
raft::host_span<size_t const>(displacements.data(), displacements.size()),
handle_.get_stream());
source_indices = std::move(allgathered_indices);

std::tie(edge_offsets, src, dst, wgt) =
Expand Down
13 changes: 7 additions & 6 deletions cpp/src/c_api/neighbor_sampling.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -951,12 +951,13 @@ struct neighbor_sampling_functor : public cugraph::c_api::abstract_functor {
label_to_comm_rank = rmm::device_uvector<label_t>(
displacements.back() + recvcounts.back(), handle_.get_stream());

cugraph::device_allgatherv(handle_.get_comms(),
local_label_to_comm_rank.begin(),
(*label_to_comm_rank).begin(),
recvcounts,
displacements,
handle_.get_stream());
cugraph::device_allgatherv(
handle_.get_comms(),
local_label_to_comm_rank.begin(),
(*label_to_comm_rank).begin(),
raft::host_span<size_t const>(recvcounts.data(), recvcounts.size()),
raft::host_span<size_t const>(displacements.data(), displacements.size()),
handle_.get_stream());

std::tie(start_vertices, *start_vertex_labels) =
cugraph::detail::shuffle_ext_vertex_value_pairs_to_local_gpu_by_vertex_partitioning(
Expand Down
8 changes: 5 additions & 3 deletions cpp/src/centrality/betweenness_centrality_impl.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -97,8 +97,8 @@ std::tuple<rmm::device_uvector<vertex_t>, rmm::device_uvector<edge_t>> brandes_b
//
// Predecessors could be a CSR if that's helpful for doing the backwards tracing
constexpr vertex_t invalid_distance = std::numeric_limits<vertex_t>::max();
constexpr int bucket_idx_cur{0};
constexpr int bucket_idx_next{1};
constexpr size_t bucket_idx_cur{0};
constexpr size_t bucket_idx_next{1};

rmm::device_uvector<edge_t> sigmas(graph_view.local_vertex_partition_range_size(),
handle.get_stream());
Expand Down Expand Up @@ -144,12 +144,14 @@ std::tuple<rmm::device_uvector<vertex_t>, rmm::device_uvector<edge_t>> brandes_b
brandes_e_op_t<vertex_t>{},
reduce_op::plus<vertex_t>());

auto next_frontier_bucket_indices = std::vector<size_t>{bucket_idx_next};
update_v_frontier(handle,
graph_view,
std::move(new_frontier),
std::move(new_sigma),
vertex_frontier,
std::vector<size_t>{bucket_idx_next},
raft::host_span<size_t const>(next_frontier_bucket_indices.data(),
next_frontier_bucket_indices.size()),
thrust::make_zip_iterator(distances.begin(), sigmas.begin()),
thrust::make_zip_iterator(distances.begin(), sigmas.begin()),
[hop] __device__(auto v, auto old_values, auto v_sigma) {
Expand Down
8 changes: 5 additions & 3 deletions cpp/src/components/weakly_connected_components_impl.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -572,13 +572,15 @@ void weakly_connected_components_impl(raft::handle_t const& handle,
num_edge_inserts.data()},
reduce_op::null());

auto next_frontier_bucket_indices =
GraphViewType::is_multi_gpu ? std::vector<size_t>{bucket_idx_next, bucket_idx_conflict}
: std::vector<size_t>{bucket_idx_next};
update_v_frontier(handle,
level_graph_view,
std::move(new_frontier_tagged_vertex_buffer),
vertex_frontier,
GraphViewType::is_multi_gpu
? std::vector<size_t>{bucket_idx_next, bucket_idx_conflict}
: std::vector<size_t>{bucket_idx_next},
raft::host_span<size_t const>(next_frontier_bucket_indices.data(),
next_frontier_bucket_indices.size()),
thrust::make_constant_iterator(0) /* dummy */,
thrust::make_discard_iterator() /* dummy */,
v_op_t<GraphViewType>{vertex_partition,
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/cores/core_number_impl.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -238,7 +238,7 @@ void core_number(raft::handle_t const& handle,
std::move(new_frontier_vertex_buffer),
std::move(delta_buffer),
vertex_frontier,
std::vector<size_t>{bucket_idx_next},
raft::host_span<size_t const>(&bucket_idx_next, size_t{1}),
core_numbers,
core_numbers,
[k_first,
Expand Down
5 changes: 4 additions & 1 deletion cpp/src/detail/permute_range.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -118,7 +118,10 @@ rmm::device_uvector<vertex_t> permute_range(raft::handle_t const& handle,
}

std::tie(permuted_integers, std::ignore) = cugraph::shuffle_values(
handle.get_comms(), permuted_integers.begin(), tx_value_counts, handle.get_stream());
handle.get_comms(),
permuted_integers.begin(),
raft::host_span<size_t const>(tx_value_counts.data(), tx_value_counts.size()),
handle.get_stream());
}

// permute locally
Expand Down
4 changes: 2 additions & 2 deletions cpp/src/link_prediction/similarity_impl.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -527,8 +527,8 @@ all_pairs_similarity(raft::handle_t const& handle,
thrust::make_zip_iterator(
gathered_v1.begin(), gathered_v2.begin(), gathered_score.begin()),
top_v1.size(),
rx_sizes,
rx_displs,
raft::host_span<size_t const>(rx_sizes.data(), rx_sizes.size()),
raft::host_span<size_t const>(rx_displs.data(), rx_displs.size()),
int{0},
handle.get_stream());

Expand Down
13 changes: 7 additions & 6 deletions cpp/src/lookup/lookup_src_dst_impl.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -228,12 +228,13 @@ struct lookup_container_t<edge_id_t, edge_type_t, vertex_t, value_t>::lookup_con
rmm::device_uvector<edge_type_t> rx_unique_types(rx_displacements.back() + rx_counts.back(),
handle.get_stream());

device_allgatherv(comm,
unique_types.begin(),
rx_unique_types.begin(),
rx_counts,
rx_displacements,
handle.get_stream());
device_allgatherv(
comm,
unique_types.begin(),
rx_unique_types.begin(),
raft::host_span<size_t const>(rx_counts.data(), rx_counts.size()),
raft::host_span<size_t const>(rx_displacements.data(), rx_displacements.size()),
handle.get_stream());
unique_types = std::move(rx_unique_types);

thrust::sort(handle.get_thrust_policy(), unique_types.begin(), unique_types.end());
Expand Down
129 changes: 76 additions & 53 deletions cpp/src/prims/detail/nbr_intersection.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -790,12 +790,13 @@ nbr_intersection(raft::handle_t const& handle,
rx_counts.begin(), rx_counts.end(), rx_displacements.begin(), size_t{0});
rmm::device_uvector<vertex_t> rx_unique_majors(rx_displacements.back() + rx_counts.back(),
handle.get_stream());
device_allgatherv(minor_comm,
unique_majors.begin(),
rx_unique_majors.begin(),
rx_counts,
rx_displacements,
handle.get_stream());
device_allgatherv(
minor_comm,
unique_majors.begin(),
rx_unique_majors.begin(),
raft::host_span<size_t const>(rx_counts.data(), rx_counts.size()),
raft::host_span<size_t const>(rx_displacements.data(), rx_displacements.size()),
handle.get_stream());
unique_majors = std::move(rx_unique_majors);

thrust::sort(handle.get_thrust_policy(), unique_majors.begin(), unique_majors.end());
Expand Down Expand Up @@ -849,12 +850,16 @@ nbr_intersection(raft::handle_t const& handle,
}

std::tie(rx_majors, rx_major_counts) =
shuffle_values(major_comm, unique_majors.begin(), tx_counts, handle.get_stream());
shuffle_values(major_comm,
unique_majors.begin(),
raft::host_span<size_t const>(tx_counts.data(), tx_counts.size()),
handle.get_stream());

std::vector<size_t> tmp_counts(major_comm_size, minor_comm_size);
std::tie(rx_group_counts, std::ignore) =
shuffle_values(major_comm,
d_tx_group_counts.begin(),
std::vector<size_t>(major_comm_size, minor_comm_size),
raft::host_span<size_t const>(tmp_counts.data(), tmp_counts.size()),
handle.get_stream());
}

Expand Down Expand Up @@ -1030,7 +1035,10 @@ nbr_intersection(raft::handle_t const& handle,
{
rmm::device_uvector<edge_t> local_degrees_for_unique_majors(size_t{0}, handle.get_stream());
std::tie(local_degrees_for_unique_majors, std::ignore) = shuffle_values(
major_comm, local_degrees_for_rx_majors.begin(), rx_major_counts, handle.get_stream());
major_comm,
local_degrees_for_rx_majors.begin(),
raft::host_span<size_t const>(rx_major_counts.data(), rx_major_counts.size()),
handle.get_stream());
major_nbr_offsets = rmm::device_uvector<edge_t>(local_degrees_for_unique_majors.size() + 1,
handle.get_stream());
(*major_nbr_offsets).set_element_to_zero_async(size_t{0}, handle.get_stream());
Expand All @@ -1043,14 +1051,17 @@ nbr_intersection(raft::handle_t const& handle,
}

std::tie(major_nbr_indices, std::ignore) = shuffle_values(
major_comm, local_nbrs_for_rx_majors.begin(), local_nbr_counts, handle.get_stream());
major_comm,
local_nbrs_for_rx_majors.begin(),
raft::host_span<size_t const>(local_nbr_counts.data(), local_nbr_counts.size()),
handle.get_stream());

if constexpr (!std::is_same_v<edge_property_value_t, cuda::std::nullopt_t>) {
std::tie(major_e_property_values, std::ignore) =
shuffle_values(major_comm,
local_e_property_values_for_rx_majors.begin(),
local_nbr_counts,
handle.get_stream());
std::tie(major_e_property_values, std::ignore) = shuffle_values(
major_comm,
local_e_property_values_for_rx_majors.begin(),
raft::host_span<size_t const>(local_nbr_counts.data(), local_nbr_counts.size()),
handle.get_stream());
}

major_to_idx_map_ptr = std::make_unique<kv_store_t<vertex_t, vertex_t, false>>(
Expand Down Expand Up @@ -1189,8 +1200,9 @@ nbr_intersection(raft::handle_t const& handle,
minor_comm,
get_dataframe_buffer_begin(vertex_pair_buffer) + rx_v_pair_displacements[minor_comm_rank],
get_dataframe_buffer_begin(vertex_pair_buffer),
rx_v_pair_counts,
rx_v_pair_displacements,
raft::host_span<size_t const>(rx_v_pair_counts.data(), rx_v_pair_counts.size()),
raft::host_span<size_t const>(rx_v_pair_displacements.data(),
rx_v_pair_displacements.size()),
handle.get_stream());

auto edge_partition =
Expand Down Expand Up @@ -1398,16 +1410,18 @@ nbr_intersection(raft::handle_t const& handle,
}
rmm::device_uvector<edge_t> gathered_nbr_intersection_sizes(
rx_v_pair_counts[minor_comm_rank] * minor_comm_size, handle.get_stream());
std::vector<size_t> rx_counts(minor_comm_size, rx_v_pair_counts[minor_comm_rank]);
device_multicast_sendrecv(
minor_comm,
rx_v_pair_nbr_intersection_sizes.begin(),
rx_v_pair_counts,
rx_v_pair_displacements,
ranks,
raft::host_span<size_t const>(rx_v_pair_counts.data(), rx_v_pair_counts.size()),
raft::host_span<size_t const>(rx_v_pair_displacements.data(),
rx_v_pair_displacements.size()),
raft::host_span<int const>(ranks.data(), ranks.size()),
gathered_nbr_intersection_sizes.begin(),
std::vector<size_t>(minor_comm_size, rx_v_pair_counts[minor_comm_rank]),
displacements,
ranks,
raft::host_span<size_t const>(rx_counts.data(), rx_counts.size()),
raft::host_span<size_t const>(displacements.data(), displacements.size()),
raft::host_span<int const>(ranks.data(), ranks.size()),
handle.get_stream());
rx_v_pair_nbr_intersection_sizes.resize(size_t{0}, handle.get_stream());
rx_v_pair_nbr_intersection_sizes.shrink_to_fit(handle.get_stream());
Expand Down Expand Up @@ -1491,16 +1505,19 @@ nbr_intersection(raft::handle_t const& handle,
rmm::device_uvector<vertex_t> gathered_nbr_intersection_indices(
rx_displacements.back() + gathered_nbr_intersection_index_rx_counts.back(),
handle.get_stream());
device_multicast_sendrecv(minor_comm,
rx_v_pair_nbr_intersection_indices.begin(),
rx_v_pair_nbr_intersection_index_tx_counts,
tx_displacements,
ranks,
gathered_nbr_intersection_indices.begin(),
gathered_nbr_intersection_index_rx_counts,
rx_displacements,
ranks,
handle.get_stream());
device_multicast_sendrecv(
minor_comm,
rx_v_pair_nbr_intersection_indices.begin(),
raft::host_span<size_t const>(rx_v_pair_nbr_intersection_index_tx_counts.data(),
rx_v_pair_nbr_intersection_index_tx_counts.size()),
raft::host_span<size_t const>(tx_displacements.data(), tx_displacements.size()),
raft::host_span<int const>(ranks.data(), ranks.size()),
gathered_nbr_intersection_indices.begin(),
raft::host_span<size_t const>(gathered_nbr_intersection_index_rx_counts.data(),
gathered_nbr_intersection_index_rx_counts.size()),
raft::host_span<size_t const>(rx_displacements.data(), rx_displacements.size()),
raft::host_span<int const>(ranks.data(), ranks.size()),
handle.get_stream());
rx_v_pair_nbr_intersection_indices.resize(size_t{0}, handle.get_stream());
rx_v_pair_nbr_intersection_indices.shrink_to_fit(handle.get_stream());

Expand All @@ -1518,32 +1535,38 @@ nbr_intersection(raft::handle_t const& handle,
handle.get_stream());

if constexpr (!std::is_same_v<edge_property_value_t, cuda::std::nullopt_t>) {
device_multicast_sendrecv(minor_comm,
rx_v_pair_nbr_intersection_e_property_values0.begin(),
rx_v_pair_nbr_intersection_index_tx_counts,
tx_displacements,
ranks,
gathered_nbr_intersection_e_property_values0.begin(),
gathered_nbr_intersection_index_rx_counts,
rx_displacements,
ranks,
handle.get_stream());
device_multicast_sendrecv(
minor_comm,
rx_v_pair_nbr_intersection_e_property_values0.begin(),
raft::host_span<size_t const>(rx_v_pair_nbr_intersection_index_tx_counts.data(),
rx_v_pair_nbr_intersection_index_tx_counts.size()),
raft::host_span<size_t const>(tx_displacements.data(), tx_displacements.size()),
raft::host_span<int const>(ranks.data(), ranks.size()),
gathered_nbr_intersection_e_property_values0.begin(),
raft::host_span<size_t const>(gathered_nbr_intersection_index_rx_counts.data(),
gathered_nbr_intersection_index_rx_counts.size()),
raft::host_span<size_t const>(rx_displacements.data(), rx_displacements.size()),
raft::host_span<int const>(ranks.data(), ranks.size()),
handle.get_stream());
rx_v_pair_nbr_intersection_e_property_values0.resize(size_t{0}, handle.get_stream());
rx_v_pair_nbr_intersection_e_property_values0.shrink_to_fit(handle.get_stream());

combined_nbr_intersection_e_property_values0.resize(
gathered_nbr_intersection_e_property_values0.size(), handle.get_stream());

device_multicast_sendrecv(minor_comm,
rx_v_pair_nbr_intersection_e_property_values1.begin(),
rx_v_pair_nbr_intersection_index_tx_counts,
tx_displacements,
ranks,
gathered_nbr_intersection_e_property_values1.begin(),
gathered_nbr_intersection_index_rx_counts,
rx_displacements,
ranks,
handle.get_stream());
device_multicast_sendrecv(
minor_comm,
rx_v_pair_nbr_intersection_e_property_values1.begin(),
raft::host_span<size_t const>(rx_v_pair_nbr_intersection_index_tx_counts.data(),
rx_v_pair_nbr_intersection_index_tx_counts.size()),
raft::host_span<size_t const>(tx_displacements.data(), tx_displacements.size()),
raft::host_span<int const>(ranks.data(), ranks.size()),
gathered_nbr_intersection_e_property_values1.begin(),
raft::host_span<size_t const>(gathered_nbr_intersection_index_rx_counts.data(),
gathered_nbr_intersection_index_rx_counts.size()),
raft::host_span<size_t const>(rx_displacements.data(), rx_displacements.size()),
raft::host_span<int const>(ranks.data(), ranks.size()),
handle.get_stream());
rx_v_pair_nbr_intersection_e_property_values1.resize(size_t{0}, handle.get_stream());
rx_v_pair_nbr_intersection_e_property_values1.shrink_to_fit(handle.get_stream());
combined_nbr_intersection_e_property_values1.resize(
Expand Down
Loading
Loading